Merging decuda_to_ptxplus with cuobjdump_to_ptxplus

Major rework for the Makefile
Deleting decuda_to_ptxplus
Removing all dependencies on boost and removing it from the dependencies of the simulator

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

26 files changed, 774 insertions, 4320 deletions

diff --git a/README b/README
index b563146..df81554 100644
--- a/README
+++ b/README
@@ -82,7 +82,6 @@ GPGPU-Sim dependencies:
 * bison
 * flex
 * zlib
-* libboost
 * CUDA Toolkit
 	
 GPGPU-Sim documentation dependencies:
@@ -104,7 +103,7 @@ dependencies besides the CUDA Toolkit.
 
 GPGPU-Sim dependencies:
 "sudo apt-get install build-essential xutils-dev bison zlib1g-dev flex
-libboost-all-dev libglu1-mesa-dev"
+libglu1-mesa-dev"
 
 GPGPU-Sim documentation dependencies:
 "sudo apt-get install doxygen graphviz"
diff --git a/cuobjdump_to_ptxplus/Makefile b/cuobjdump_to_ptxplus/Makefile
index 2f3c202..7f58f69 100644
--- a/cuobjdump_to_ptxplus/Makefile
+++ b/cuobjdump_to_ptxplus/Makefile
@@ -1,11 +1,9 @@
-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
+CCFLAGS		= -ggdb -fPIC
 CXX			= g++
 CXXFLAGS	= ${CCFLAGS}
 LD			= g++
@@ -14,71 +12,18 @@ LDFLAGS		= ${CCFLAGS}
 
 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
+HEADER_PARSER_OBJECTS = header_parser.o header_lexer.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)
-	${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
+MAKEFLAGS += --no-builtin-rules
 
-headerParser.o : headerParser.cpp ../decuda_to_ptxplus/stringList.h ../decuda_to_ptxplus/decudaInst.h ./decudaInstList.h
-	${CXX} ${CXXFLAGS} -c -Igenerated -oheaderParser.o headerParser.cpp
+.SUFFIXES:
+.SECONDARY:
 
-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
+cuobjdump_to_ptxplus: stringList.o cuobjdumpInst.o cuobjdumpInstList.o cuobjdump_to_ptxplus.o $(PTX_PARSER_OBJECTS) $(SASS_PARSER_OBJECTS) $(ELF_PARSER_OBJECTS) $(HEADER_PARSER_OBJECTS)
+	${LD} ${LDFLAGS} -o $@ stringList.o cuobjdumpInst.o cuobjdumpInstList.o cuobjdump_to_ptxplus.o $(PTX_PARSER_OBJECTS) $(SASS_PARSER_OBJECTS) $(ELF_PARSER_OBJECTS) $(HEADER_PARSER_OBJECTS)
 
 lex.ptx_.c : ../src/cuda-sim/ptx.l
 	${LEX} ${LEXFLAGS} ../src/cuda-sim/ptx.l
@@ -97,37 +42,27 @@ lex.ptx_.o : lex.ptx_.c ptx.tab.h
 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
+%_lexer.cc: %.l %_parser.hh
 	$(LEX) $(LEXFLAGS) -P $*_ -o$@ $<
 
-%_lexer.cpp: %.l %_parser.hpp
-	$(LEX) $(LEXFLAGS) -P $*_ -o$@ $<
-
-%_parser.c: %.y
-	$(YACC) $(YFLAGS) -p $*_ -o$@ $<
-
-%_parser.cpp: %.y
+%_parser.cc: %.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.hh: %_parser.cc
 	:
 
-%_parser.hpp: %_parser.cpp
-	:
+%.o: %.cc
+	$(CXX) ${CXXFLAGS} -c -o $@ $<
 
-cuobjdump_parser.o: cuobjdump_parser.c
-	${CC} ${CCFLAGS} -c -o $@ $<
+%.o: %.cpp %.h
+	$(CXX) ${CXXFLAGS} -c -o $@ $<
 
-cuobjdump_lexer.o: cuobjdump_lexer.c
+%.o: %.c %.h
 	${CC} ${CCFLAGS} -c -o $@ $<
 
-%.o: %.cpp
-	$(CXX) ${CXXFLAGS} -c -o $@ $<
-
 SRCS = $(shell ls *.cc)
 
 depend:
@@ -136,22 +71,53 @@ depend:
 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 ./header_parser.output ./header_parser.o header_parser.cc header_parser.hh header_lexer.cc header_lexer.hh
+	rm -f ./stringList.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.hpp ./sass_parser.output ./sass_parser.h
 	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
+	rm -f ./header_lexer.o header_lexer.c header_lexer.h sass_lexer.c sass_lexer.h sass_parser.c
+	rm -f ./sass_lexer.cc ./sass_parser.cc ./sass_parser.hh ./sass_lexer.hh
 # DO NOT DELETE
 
-cuobjdumpInst.o: cuobjdumpInst.h
+cuobjdumpInst.o: cuobjdumpInst.h /usr/include/string.h
+cuobjdumpInst.o: /usr/include/features.h /usr/include/sys/cdefs.h
+cuobjdumpInst.o: /usr/include/bits/wordsize.h /usr/include/gnu/stubs.h
+cuobjdumpInst.o: /usr/include/gnu/stubs-64.h /usr/include/xlocale.h
+cuobjdumpInst.o: /usr/include/stdio.h /usr/include/bits/types.h
+cuobjdumpInst.o: /usr/include/bits/typesizes.h /usr/include/libio.h
+cuobjdumpInst.o: /usr/include/_G_config.h /usr/include/wchar.h
+cuobjdumpInst.o: /usr/include/bits/stdio_lim.h
+cuobjdumpInst.o: /usr/include/bits/sys_errlist.h /usr/include/stdlib.h
+cuobjdumpInst.o: /usr/include/sys/types.h /usr/include/time.h
+cuobjdumpInst.o: /usr/include/endian.h /usr/include/bits/endian.h
+cuobjdumpInst.o: /usr/include/bits/byteswap.h /usr/include/sys/select.h
+cuobjdumpInst.o: /usr/include/bits/select.h /usr/include/bits/sigset.h
+cuobjdumpInst.o: /usr/include/bits/time.h /usr/include/sys/sysmacros.h
+cuobjdumpInst.o: /usr/include/bits/pthreadtypes.h /usr/include/alloca.h
+cuobjdumpInst.o: stringList.h
 cuobjdumpInstList.o: cuobjdumpInstList.h cuobjdumpInst.h
-cuobjdump_to_ptxplus.o: cuobjdumpInstList.h cuobjdumpInst.h
+cuobjdumpInstList.o: /usr/include/string.h /usr/include/features.h
+cuobjdumpInstList.o: /usr/include/sys/cdefs.h /usr/include/bits/wordsize.h
+cuobjdumpInstList.o: /usr/include/gnu/stubs.h /usr/include/gnu/stubs-64.h
+cuobjdumpInstList.o: /usr/include/xlocale.h /usr/include/stdio.h
+cuobjdumpInstList.o: /usr/include/bits/types.h /usr/include/bits/typesizes.h
+cuobjdumpInstList.o: /usr/include/libio.h /usr/include/_G_config.h
+cuobjdumpInstList.o: /usr/include/wchar.h /usr/include/bits/stdio_lim.h
+cuobjdumpInstList.o: /usr/include/bits/sys_errlist.h /usr/include/stdlib.h
+cuobjdumpInstList.o: /usr/include/sys/types.h /usr/include/time.h
+cuobjdumpInstList.o: /usr/include/endian.h /usr/include/bits/endian.h
+cuobjdumpInstList.o: /usr/include/bits/byteswap.h /usr/include/sys/select.h
+cuobjdumpInstList.o: /usr/include/bits/select.h /usr/include/bits/sigset.h
+cuobjdumpInstList.o: /usr/include/bits/time.h /usr/include/sys/sysmacros.h
+cuobjdumpInstList.o: /usr/include/bits/pthreadtypes.h /usr/include/alloca.h
+cuobjdumpInstList.o: stringList.h
 cuobjdump_to_ptxplus.o: /usr/include/stdio.h /usr/include/features.h
 cuobjdump_to_ptxplus.o: /usr/include/sys/cdefs.h /usr/include/bits/wordsize.h
 cuobjdump_to_ptxplus.o: /usr/include/gnu/stubs.h /usr/include/gnu/stubs-64.h
@@ -159,4 +125,56 @@ cuobjdump_to_ptxplus.o: /usr/include/bits/types.h
 cuobjdump_to_ptxplus.o: /usr/include/bits/typesizes.h /usr/include/libio.h
 cuobjdump_to_ptxplus.o: /usr/include/_G_config.h /usr/include/wchar.h
 cuobjdump_to_ptxplus.o: /usr/include/bits/stdio_lim.h
-cuobjdump_to_ptxplus.o: /usr/include/bits/sys_errlist.h
+cuobjdump_to_ptxplus.o: /usr/include/bits/sys_errlist.h cuobjdumpInstList.h
+cuobjdump_to_ptxplus.o: cuobjdumpInst.h /usr/include/string.h
+cuobjdump_to_ptxplus.o: /usr/include/xlocale.h /usr/include/stdlib.h
+cuobjdump_to_ptxplus.o: /usr/include/sys/types.h /usr/include/time.h
+cuobjdump_to_ptxplus.o: /usr/include/endian.h /usr/include/bits/endian.h
+cuobjdump_to_ptxplus.o: /usr/include/bits/byteswap.h
+cuobjdump_to_ptxplus.o: /usr/include/sys/select.h /usr/include/bits/select.h
+cuobjdump_to_ptxplus.o: /usr/include/bits/sigset.h /usr/include/bits/time.h
+cuobjdump_to_ptxplus.o: /usr/include/sys/sysmacros.h
+cuobjdump_to_ptxplus.o: /usr/include/bits/pthreadtypes.h
+cuobjdump_to_ptxplus.o: /usr/include/alloca.h stringList.h
+elf_lexer.o: /usr/include/stdio.h /usr/include/features.h
+elf_lexer.o: /usr/include/sys/cdefs.h /usr/include/bits/wordsize.h
+elf_lexer.o: /usr/include/gnu/stubs.h /usr/include/gnu/stubs-64.h
+elf_lexer.o: /usr/include/bits/types.h /usr/include/bits/typesizes.h
+elf_lexer.o: /usr/include/libio.h /usr/include/_G_config.h
+elf_lexer.o: /usr/include/wchar.h /usr/include/bits/stdio_lim.h
+elf_lexer.o: /usr/include/bits/sys_errlist.h /usr/include/string.h
+elf_lexer.o: /usr/include/xlocale.h /usr/include/errno.h
+elf_lexer.o: /usr/include/bits/errno.h /usr/include/linux/errno.h
+elf_lexer.o: /usr/include/asm/errno.h /usr/include/asm-generic/errno.h
+elf_lexer.o: /usr/include/asm-generic/errno-base.h /usr/include/stdlib.h
+elf_lexer.o: /usr/include/sys/types.h /usr/include/time.h
+elf_lexer.o: /usr/include/endian.h /usr/include/bits/endian.h
+elf_lexer.o: /usr/include/bits/byteswap.h /usr/include/sys/select.h
+elf_lexer.o: /usr/include/bits/select.h /usr/include/bits/sigset.h
+elf_lexer.o: /usr/include/bits/time.h /usr/include/sys/sysmacros.h
+elf_lexer.o: /usr/include/bits/pthreadtypes.h /usr/include/alloca.h
+elf_lexer.o: elf_parser.hh /usr/include/unistd.h
+elf_lexer.o: /usr/include/bits/posix_opt.h /usr/include/bits/confname.h
+elf_lexer.o: /usr/include/getopt.h
+elf_parser.o: /usr/include/stdio.h /usr/include/features.h
+elf_parser.o: /usr/include/sys/cdefs.h /usr/include/bits/wordsize.h
+elf_parser.o: /usr/include/gnu/stubs.h /usr/include/gnu/stubs-64.h
+elf_parser.o: /usr/include/bits/types.h /usr/include/bits/typesizes.h
+elf_parser.o: /usr/include/libio.h /usr/include/_G_config.h
+elf_parser.o: /usr/include/wchar.h /usr/include/bits/stdio_lim.h
+elf_parser.o: /usr/include/bits/sys_errlist.h cuobjdumpInstList.h
+elf_parser.o: cuobjdumpInst.h /usr/include/string.h /usr/include/xlocale.h
+elf_parser.o: /usr/include/stdlib.h /usr/include/sys/types.h
+elf_parser.o: /usr/include/time.h /usr/include/endian.h
+elf_parser.o: /usr/include/bits/endian.h /usr/include/bits/byteswap.h
+elf_parser.o: /usr/include/sys/select.h /usr/include/bits/select.h
+elf_parser.o: /usr/include/bits/sigset.h /usr/include/bits/time.h
+elf_parser.o: /usr/include/sys/sysmacros.h /usr/include/bits/pthreadtypes.h
+elf_parser.o: /usr/include/alloca.h stringList.h
+stringList.o: stringList.h /usr/include/stdio.h /usr/include/features.h
+stringList.o: /usr/include/sys/cdefs.h /usr/include/bits/wordsize.h
+stringList.o: /usr/include/gnu/stubs.h /usr/include/gnu/stubs-64.h
+stringList.o: /usr/include/bits/types.h /usr/include/bits/typesizes.h
+stringList.o: /usr/include/libio.h /usr/include/_G_config.h
+stringList.o: /usr/include/wchar.h /usr/include/bits/stdio_lim.h
+stringList.o: /usr/include/bits/sys_errlist.h
diff --git a/cuobjdump_to_ptxplus/cuobjdumpInst.cc b/cuobjdump_to_ptxplus/cuobjdumpInst.cc
index 524c0da..e0d5557 100644
--- a/cuobjdump_to_ptxplus/cuobjdumpInst.cc
+++ b/cuobjdump_to_ptxplus/cuobjdumpInst.cc
@@ -1,6 +1,10 @@
-#include "cuobjdumpInst.h"
+// External includes
 #include <sstream>
 #include <iostream>
+#include <cassert>
+
+// Local includes
+#include "cuobjdumpInst.h"
 
 extern void output(const char * text);
 
@@ -8,6 +12,191 @@ extern void output(const char * text);
 cuobjdumpInst::cuobjdumpInst()
 {
 	//initilize everything to empty
+	m_label = "";
+	m_predicate = new stringList();
+	m_base = "";
+	m_baseModifiers = new stringList();
+	m_typeModifiers = new stringList();
+	m_operands = new stringList();
+	m_predicateModifiers = new stringList();
+}
+
+void cuobjdumpInst::printCuobjdumpInst()
+{
+	/*TODO: print label here*/
+	/*TODO: print predicate here*/
+	/*std::cout << "Instruction Base: " << m_base << "\n";
+
+	std::cout << "Instruction Modifiers: ";
+	m_baseModifiers->printStringList();
+	std::cout << "\n";
+
+	std::cout << "Operand types: ";
+	m_typeModifiers->printStringList();
+	std::cout << "\n";
+
+	std::cout << "Operands: ";
+	m_operands->printStringList();
+	std::cout << "\n\n";*/
+
+	std::cout << m_base << " ";
+	m_baseModifiers->printStringList();
+	std::cout << " ";
+	m_typeModifiers->printStringList();
+	std::cout << " ";
+	m_operands->printStringList();
+	std::cout << "\n";
+}
+
+// Just prints the base and operands
+void cuobjdumpInst::printHeaderPtx()
+{
+	output(m_base); output(" ");
+
+	stringListPiece* currentPiece;
+
+	currentPiece = m_baseModifiers->getListStart();
+	for(int i=0; (i<m_baseModifiers->getSize())&&(currentPiece!=NULL); i++)
+	{
+		output(" "); output(currentPiece->stringText);
+		currentPiece = currentPiece->nextString;
+	}
+
+	currentPiece = m_operands->getListStart();
+	for(int i=0; (i<m_operands->getSize())&&(currentPiece!=NULL); i++)
+	{
+		output(" "); output(currentPiece->stringText);
+		currentPiece = currentPiece->nextString;
+	}
+}
+
+//retreive instruction mnemonic
+const char* cuobjdumpInst::getBase()
+{
+	return m_base;
+}
+
+stringList* cuobjdumpInst::getTypeModifiers()
+{
+        return m_typeModifiers;
+}
+
+//print out .version and .target header lines
+bool cuobjdumpInst::printHeaderInst()
+{
+	if(strcmp(m_base, ".version")==0)
+	{
+		output(m_base); output(" ");
+
+		stringListPiece* currentPiece = m_operands->getListStart();
+		output(currentPiece->stringText);
+		currentPiece = currentPiece->nextString;
+
+		if(currentPiece!=NULL)
+		{
+			output("."); output(currentPiece->stringText);
+		}
+      output("+");
+		output("\n");
+	}
+	else if(strcmp(m_base, ".target")==0)
+	{
+		output(m_base); output(" ");
+
+		stringListPiece* currentPiece = m_operands->getListStart();
+		output(currentPiece->stringText);
+		currentPiece = currentPiece->nextString;
+
+		while(currentPiece!=NULL)
+		{
+			output(", "); output(currentPiece->stringText);
+			currentPiece = currentPiece->nextString;
+		}
+		output("\n");
+	}
+	else if(strcmp(m_base, ".tex")==0)
+	{
+		output(m_base); output(" ");
+
+		stringListPiece* currentPiece;
+
+
+		currentPiece = m_baseModifiers->getListStart();
+		output(currentPiece->stringText); output(" ");
+		currentPiece = currentPiece->nextString;
+
+		while(currentPiece!=NULL)
+		{
+			output(" "); output(currentPiece->stringText);
+			currentPiece = currentPiece->nextString;
+		}
+
+		currentPiece = m_operands->getListStart();
+		output(currentPiece->stringText);
+		currentPiece = currentPiece->nextString;
+
+		while(currentPiece!=NULL)
+		{
+			output(" "); output(currentPiece->stringText);
+			currentPiece = currentPiece->nextString;
+		}
+		output(";\n");
+	}
+	else
+	{
+		return false;
+	}
+	return true;
+}
+
+void cuobjdumpInst::setBase(const char* setBaseValue)
+{
+	m_base = setBaseValue;
+}
+
+void cuobjdumpInst::addBaseModifier(const char* addBaseMod)
+{
+	stringListPiece* tempPiece = new stringListPiece;
+	tempPiece->stringText = addBaseMod;
+
+	m_baseModifiers->add(tempPiece);
+}
+
+void cuobjdumpInst::addTypeModifier(const char* addTypeMod)
+{
+	stringListPiece* tempPiece = new stringListPiece;
+	tempPiece->stringText = addTypeMod;
+
+	m_typeModifiers->add(tempPiece);
+}
+
+void cuobjdumpInst::addOperand(const char* addOp)
+{
+	stringListPiece* tempPiece = new stringListPiece;
+	tempPiece->stringText = addOp;
+
+	m_operands->add(tempPiece);
+}
+
+void cuobjdumpInst::setPredicate(const char* setPredicateValue)
+{
+	stringListPiece* tempPiece = new stringListPiece;
+	tempPiece->stringText = setPredicateValue;
+
+	m_predicate->add(tempPiece);
+}
+
+void cuobjdumpInst::addPredicateModifier(const char* addPredicateMod)
+{
+	stringListPiece* tempPiece = new stringListPiece;
+	tempPiece->stringText = addPredicateMod;
+
+	m_predicateModifiers->add(tempPiece);
+}
+
+void cuobjdumpInst::setLabel(const char* setLabelValue)
+{
+	m_label = setLabelValue;
 }
 
 bool cuobjdumpInst::checkCubojdumpLabel(std::list<std::string> labelList, std::string label)
@@ -182,7 +371,7 @@ void cuobjdumpInst::printCuobjdumpBaseModifiers()
 		}
 		else if( strcmp(modString, ".s")==0 )
 		{
-			//".s" is the same as ".join" in decuda.
+			//".s" is the same as ".join" in cuobjdump.
 		}
 		else if( strcmp(modString, ".sfu")==0 )
 		{
diff --git a/cuobjdump_to_ptxplus/cuobjdumpInst.h b/cuobjdump_to_ptxplus/cuobjdumpInst.h
index 20d31be..8cf0ebe 100644
--- a/cuobjdump_to_ptxplus/cuobjdumpInst.h
+++ b/cuobjdump_to_ptxplus/cuobjdumpInst.h
@@ -1,12 +1,45 @@
-#include "decudaInstList.h"
+#ifndef _CUOBJDUMPINST_H_
+#define _CUOBJDUMPINST_H_
 
-class cuobjdumpInst: public decudaInst
+// External includes
+#include <string.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <list>
+
+// Local includes
+//#include "cuobjdumpInstList.h"
+#include "stringList.h"
+
+class cuobjdumpInst
 {
+protected:
+	//instruction data
+	const char* m_label; //instruction label
+	stringList* m_predicate; //instruction predicate
+	const char* m_base; //instruction mnemonic
+	stringList* m_baseModifiers; //base modifiers
+	stringList* m_typeModifiers; //operand types
+	stringList* m_operands; //operands
+	stringList* m_predicateModifiers; //predicate modifiers
 
 public:
 	//Constructor
 	cuobjdumpInst();
 
+	//accessors
+	const char* getBase();
+	stringList* getTypeModifiers();
+
+	//Mutators
+	void setLabel(const char* setLabelValue);
+	void setPredicate(const char* setPredicateValue);
+	void addPredicateModifier(const char* addPredicateMod);
+	void setBase(const char* setBaseValue);
+	void addBaseModifier(const char* addBaseMod);
+	void addTypeModifier(const char* addTypeMod);
+	void addOperand(const char* addOp);
+
 	bool checkCubojdumpLabel(std::list<std::string> labelList, std::string label);
 
 	void printCuobjdumpLabel(std::list<std::string> labelList);
@@ -18,5 +51,11 @@ public:
 	void printCuobjdumpOperands();
 
 	void printCuobjdumpPtxPlus(std::list<std::string> labelList, std::list<std::string> texList);
+
+	//print representation
+	bool printHeaderInst();
+	void printCuobjdumpInst();
+	void printHeaderPtx();
 };
 
+#endif //_CUOBJDUMPINST_H_
diff --git a/cuobjdump_to_ptxplus/cuobjdumpInstList.cc b/cuobjdump_to_ptxplus/cuobjdumpInstList.cc
index 5eaf67f..45c8f9d 100644
--- a/cuobjdump_to_ptxplus/cuobjdumpInstList.cc
+++ b/cuobjdump_to_ptxplus/cuobjdumpInstList.cc
@@ -1,6 +1,17 @@
+#include <sstream>
 #include <iostream>
+#include <cassert>
 #include "cuobjdumpInstList.h"
 
+#define P_DEBUG 1
+#define DPRINTF(...) \
+   if(P_DEBUG) { \
+      printf("(%s:%u) ", __FILE__, __LINE__); \
+      printf(__VA_ARGS__); \
+      printf("\n"); \
+      fflush(stdout); \
+   }
+
 extern void output(const char * text);
 
 //Constructor
@@ -11,6 +22,289 @@ cuobjdumpInstList::cuobjdumpInstList()
 
 
 
+// add to tex list
+void cuobjdumpInstList::addTex(std::string tex)
+{
+	std::string origTex = tex;
+	DPRINTF("cuobjdumpInstList::addTex tex=%s", tex.c_str());
+	// If $tex# tex from cuobjdump, then use index to get real tex name
+	if(tex.substr(0, 4) == "$tex") {
+		tex = tex.substr(4, tex.size()-4);
+		int texNum = atoi(tex.c_str());
+		if(texNum >= m_realTexList.size()) {
+			output("ERROR: tex does not exist in real tex list from ptx.\n.");
+			assert(0);
+		}
+
+		std::list<std::string>::iterator itex = m_realTexList.begin();
+		for(int i=0; i<texNum; i++) itex++;
+		origTex = *itex;
+	}
+	// Otherwise, tex from original ptx
+	else {
+		m_realTexList.push_back(tex);
+	}
+
+	// Add the tex to instruction operand list
+	//char* texName = new char [strlen(origTex.c_str())+1];
+	//strcpy(texName, origTex.c_str());
+	//getListEnd().addOperand(texName);
+}
+
+void cuobjdumpInstList::setLastEntryName(std::string entryName)
+{
+	m_entryList.back().m_entryName = entryName;
+}
+
+// create new global constant memory "bank"
+void cuobjdumpInstList::addConstMemory(int index)
+{
+	constMemory newConstMem;
+	newConstMem.index = index;
+	newConstMem.entryIndex = 0;
+	m_constMemoryList.push_back(newConstMem);
+}
+
+//add cuobjdumpInst to the last entry in entry list
+int cuobjdumpInstList::add(cuobjdumpInst* newCuobjdumpInst)
+{
+	if(m_entryList.size() == 0) {
+		//output("ERROR: Adding an instruction before entry.\n");
+		addEntry("");
+		//assert(0);
+	}
+
+	m_entryList.back().m_instList.push_back(*newCuobjdumpInst);
+
+	return m_entryList.size();
+}
+
+// add a new entry
+int cuobjdumpInstList::addEntry(std::string entryName)
+{
+	cuobjdumpEntry newEntry;
+	newEntry.m_largestRegIndex = -1;
+	newEntry.m_largestOfsRegIndex = -1;
+	newEntry.m_largestPredIndex = -1;
+	newEntry.m_reg124 = false;
+	newEntry.m_oreg127 = false;
+	newEntry.m_lMemSize = -1;
+
+	newEntry.m_entryName = entryName;
+
+
+   // Fill opPerCycle histogram with values
+   newEntry.m_opPerCycleHistogram.insert( std::pair<std::string,int>("OP_1", 0) );
+   newEntry.m_opPerCycleHistogram.insert( std::pair<std::string,int>("OP_2", 0) );
+   newEntry.m_opPerCycleHistogram.insert( std::pair<std::string,int>("OP_8", 0) );
+
+
+	m_entryList.push_back(newEntry);
+	return m_entryList.size();
+}
+
+// print out .version and .target headers
+void cuobjdumpInstList::printHeaderInstList()
+{
+	// These should be in the first entry
+	cuobjdumpEntry e_first = m_entryList.front();
+
+	std::list<cuobjdumpInst>::iterator currentInst;
+	for(currentInst=e_first.m_instList.begin(); currentInst!=e_first.m_instList.end(); ++currentInst)
+	{
+		if(!(currentInst->printHeaderInst()))
+		{
+			break;
+		}
+	}
+	for (	std::list<std::string>::iterator iter = m_realTexList.begin();
+			iter != m_realTexList.end();
+			iter ++) {
+		output(".tex .u64 ");
+		output((*iter).c_str());
+		output(";\n");
+	}
+}
+
+bool cuobjdumpInstList::findEntry(std::string entryName, cuobjdumpEntry& entry) {
+	std::list<cuobjdumpEntry>::iterator e;
+
+	std::string entryNameS = entryName;
+
+	for(e=m_entryList.begin(); e!=m_entryList.end(); ++e) {
+		if( e->m_entryName == entryNameS) {
+			entry = *e;
+			return true;
+		}
+	}
+
+	return false;
+}
+
+// get the list of real tex names
+std::list<std::string> cuobjdumpInstList::getRealTexList() {
+	return m_realTexList;
+}
+
+// set the list of real tex names
+void cuobjdumpInstList::setRealTexList(std::list<std::string> realTexList) {
+	m_realTexList = realTexList;
+}
+
+// add value to const memory
+void cuobjdumpInstList::addConstMemoryValue(std::string constMemoryValue)
+{
+	m_constMemoryList.back().m_constMemory.push_back(constMemoryValue);
+}
+
+void cuobjdumpInstList::addConstMemoryValue2(std::string constMemoryValue)
+{
+	m_constMemoryList2.back().m_constMemory.push_back(constMemoryValue);
+}
+
+// set type of constant memory
+void cuobjdumpInstList::setConstMemoryType(const char* type)
+{
+	m_constMemoryList.back().type = type;
+}
+
+void cuobjdumpInstList::setConstMemoryType2(const char* type)
+{
+	m_constMemoryList2.back().type = type;
+}
+
+//retrieve point to list end
+cuobjdumpInst cuobjdumpInstList::getListEnd()
+{
+	return m_entryList.back().m_instList.back();
+}
+
+// print out predicate names
+void cuobjdumpInstList::printPredNames(cuobjdumpEntry entry)
+{
+	if( entry.m_largestPredIndex >= 0) {
+		char out[30];
+		sprintf(out, "\t.reg .pred $p<%d>;", entry.m_largestPredIndex+1);
+		output(out);
+		output("\n");
+	}
+
+}
+
+// print reg124 and set its value to 0
+void cuobjdumpInstList::printOutOfBoundRegisters(cuobjdumpEntry entry)
+{
+	if( entry.m_reg124 == true ) {
+		output("\n");
+		output("\t.reg .u32 $r124;\n");
+		output("\tmov.u32 $r124, 0x00000000;\n");
+	}
+	if( entry.m_oreg127 == true) {
+		output("\n");
+		output("\t.reg .u32 $o127;\n");
+	}
+}
+
+// print out register names
+void cuobjdumpInstList::printRegNames(cuobjdumpEntry entry)
+{
+	if( entry.m_largestRegIndex >= 0) {
+		char out[30];
+		sprintf(out, "\t.reg .u32 $r<%d>;", entry.m_largestRegIndex+1);
+		output(out);
+		output("\n");
+	}
+
+	if( entry.m_largestOfsRegIndex >= 0) {
+		char out[30];
+		sprintf(out, "\t.reg .u32 $ofs<%d>;", entry.m_largestOfsRegIndex+1);
+		output(out);
+		output("\n");
+	}
+}
+
+// print const memory directive
+void cuobjdumpInstList::printMemory()
+{
+
+	// Constant memory
+
+	for(std::list<constMemory>::iterator i=m_constMemoryList.begin(); i!=m_constMemoryList.end(); ++i) {
+		char line[40];
+
+		// Global or entry specific
+		if(i->entryIndex == 0)
+			sprintf(line, ".const %s constant0[%d] = {", i->type, i->m_constMemory.size());
+		else
+			sprintf(line, ".const %s ce%dc%d[%d] = {", i->type, i->entryIndex, i->index, i->m_constMemory.size());
+
+		output(line);
+
+		std::list<std::string>::iterator j;
+		int l=0;
+		for(j=i->m_constMemory.begin(); j!=i->m_constMemory.end(); ++j) {
+			if(j!=i->m_constMemory.begin())
+				output(", ");
+			if( (l++ % 4) == 0) output("\n          ");
+			output(j->c_str());
+		}
+		output("\n};\n\n");
+	}
+
+
+	for(std::list<constMemory2>::iterator i=m_constMemoryList2.begin(); i!=m_constMemoryList2.end(); ++i) {
+		char line[1024];
+
+		// Global or entry specific
+		sprintf(line, ".const %s constant1%s[%d] = {", i->type, i->kernel, i->m_constMemory.size());
+
+		output(line);
+
+		std::list<std::string>::iterator j;
+		int l=0;
+		for(j=i->m_constMemory.begin(); j!=i->m_constMemory.end(); ++j) {
+			if(j!=i->m_constMemory.begin())
+				output(", ");
+			if( (l++ % 4) == 0) output("\n          ");
+			output(j->c_str());
+		}
+		output("\n};\n\n");
+	}
+
+	// Next, print out the local memory declaration
+	std::list<cuobjdumpEntry>::iterator e;
+	int eIndex=1; // entry index starts from 1 from the first blank entry is missing here (only in header entry list)
+	for(e=m_entryList.begin(); e!=m_entryList.end(); ++e) {
+		if(e->m_lMemSize > 0) {
+			std::stringstream ssout;
+			ssout << ".local .b8 l" << eIndex << "[" << e->m_lMemSize << "];" << std::endl;
+			output(ssout.str().c_str());
+		}
+		eIndex++;
+	}
+	output("\n");
+
+	// Next, print out the global memory declaration
+	std::list<globalMemory>::iterator g;
+	for(g=m_globalMemoryList.begin(); g!=m_globalMemoryList.end(); ++g) {
+		std::stringstream out;
+		out << ".global .b8 " << g->name << "[" << g->bytes << "];" << std::endl;
+		output(out.str().c_str());
+	}
+	output("\n");
+
+	// Next, print out constant memory pointers
+	std::list<constMemoryPtr>::iterator cp;
+	for(cp=m_constMemoryPtrList.begin(); cp!=m_constMemoryPtrList.end(); ++cp) {
+		std::stringstream out;
+		out << ".const .b8 " << cp->name << "[" << cp->bytes << "];" << std::endl;
+		out << ".constptr " << cp->name << ", " << cp->destination << ", " << cp->offset << ";" << std::endl;
+		output(out.str().c_str());
+	}
+	output("\n");
+
+}
+
 
 //TODO: Some register processing work is supposed to be done here.
 void cuobjdumpInstList::addCuobjdumpRegister(std::string reg, bool lo)
@@ -303,164 +597,25 @@ void cuobjdumpInstList::addEntryLocalMemory(int value, int 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;
+	std::list<cuobjdumpEntry>::iterator e;
 	for(e=m_entryList.begin(); e!=m_entryList.end(); ++e) {
 
 
 
 
-		for(	std::list<decudaInst>::iterator currentInst=e->m_instList.begin();
+		for(	std::list<cuobjdumpInst>::iterator currentInst=e->m_instList.begin();
 				currentInst!=e->m_instList.end();
 				++currentInst) {
 			// Output the instruction
 			output("\t");
-			currentInst->printDecudaInst();
+			currentInst->printCuobjdumpInst();
 			output("\n");
 		}
 	}
@@ -486,18 +641,18 @@ void cuobjdumpInstList::printCuobjdumpPtxPlusList(cuobjdumpInstList* headerInfo)
 	printMemory();
 	printCuobjdumpLocalMemory();
 	// Each entry
-	std::list<decudaEntry>::reverse_iterator e;
+	std::list<cuobjdumpEntry>::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;
+		cuobjdumpEntry headerEntry;
 
 		if( headerInfo->findEntry(e->m_entryName, headerEntry) ) {
 			// Entry for current header found, print it out
-			std::list<decudaInst>::iterator headerInstIter;
+			std::list<cuobjdumpInst>::iterator headerInstIter;
 			for(headerInstIter=headerEntry.m_instList.begin();
 				headerInstIter!=headerEntry.m_instList.end();
 				++headerInstIter) {
@@ -517,7 +672,7 @@ void cuobjdumpInstList::printCuobjdumpPtxPlusList(cuobjdumpInstList* headerInfo)
 				output("\n");
 				output("{\n");
 			} else {
-				output("Mismatch in entry names between decuda output and original ptx file.\n");
+				output("Mismatch in entry names between cuobjdump output and original ptx file.\n");
 				assert(0);
 			}
 		}
@@ -527,13 +682,12 @@ void cuobjdumpInstList::printCuobjdumpPtxPlusList(cuobjdumpInstList* headerInfo)
 		printOutOfBoundRegisters(*e);
 		output("\n");
 
-		for(std::list<decudaInst>::iterator currentInst=e->m_instList.begin(); currentInst!=e->m_instList.end(); ++currentInst){
+		for(std::list<cuobjdumpInst>::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);
+			currentInst->printCuobjdumpPtxPlus(e->m_labelList, this->m_realTexList);
 			output("\n");
 		}
 		output("\n\tl_exit: exit;\n");
diff --git a/cuobjdump_to_ptxplus/cuobjdumpInstList.h b/cuobjdump_to_ptxplus/cuobjdumpInstList.h
index 1bc8f10..7cb3a32 100644
--- a/cuobjdump_to_ptxplus/cuobjdumpInstList.h
+++ b/cuobjdump_to_ptxplus/cuobjdumpInstList.h
@@ -1,5 +1,73 @@
+#ifndef _CUOBJDUMPINSTLIST_H_
+#define _CUOBJDUMPINSTLIST_H_
+
+// External includes
+#include <list>
+#include <map>
+#include <string>
+
+// Local includes
 #include "cuobjdumpInst.h"
 
+// Used for entry specific constant memory segments (c1)
+struct constMemory
+{
+	int index;
+	int entryIndex;
+	const char* type;
+	std::list<std::string> m_constMemory;
+};
+
+struct constMemory2
+{
+	const char* kernel;
+	const char* type;
+	std::list<std::string> m_constMemory;
+};
+
+// Used for uninitialized constant memory (globally defined)
+struct constMemoryPtr
+{
+	int bytes;
+	std::string name;
+
+	std::string destination;
+	int offset;
+};
+
+// Used for global memory segments
+struct globalMemory
+{
+	int offset;
+	int bytes;
+	std::string name;
+};
+
+struct cuobjdumpEntry
+{
+	//char* m_entryName;
+	std::string m_entryName;
+	std::list<cuobjdumpInst> m_instList;	// List of cuobjdump instructions
+
+	// Register list
+	int m_largestRegIndex;
+	int m_largestOfsRegIndex;
+	bool m_reg124;
+	bool m_oreg127;
+
+	// Predicate list
+	int m_largestPredIndex;
+
+	// Local memory size
+	int m_lMemSize;
+
+	//use for recording used labels
+	std::list<std::string> m_labelList;
+
+   // Histogram for operation per cycle count
+   std::map<std::string, int> m_opPerCycleHistogram;
+};
+
 // Used for local memory segments
 struct localMemory
 {
@@ -7,22 +75,46 @@ struct localMemory
 	int entryIndex;
 };
 
-class cuobjdumpInstList: public decudaInstList
+class cuobjdumpInstList
 {
 protected:
-	std::string parseCuobjdumpPredicate(std::string pred);
+	std::list<cuobjdumpEntry> m_entryList;
+	std::list<constMemory> m_constMemoryList;
+	std::list<constMemory2> m_constMemoryList2;
+	std::list<globalMemory> m_globalMemoryList;
+
 	int m_kernelCount;
 	std::map<std::string,int>kernelcmemmap;
 	std::map<std::string,int>kernellmemmap;
 	std::list<localMemory> m_localMemoryList;
+	std::list<std::string>  m_realTexList;	// Stores the real names of tex variables
+	std::list<constMemoryPtr> m_constMemoryPtrList;
+
+	// Functions:
+	std::string parseCuobjdumpPredicate(std::string pred);
+	void printMemory();// Print const memory directives
+	// Print register names
+	void printRegNames(cuobjdumpEntry entry);
+	void printOutOfBoundRegisters(cuobjdumpEntry entry);
 
+	// Print predicate names
+	void printPredNames(cuobjdumpEntry entry);
 public:
 	//Constructor
 	cuobjdumpInstList();
 
+	cuobjdumpInst getListEnd();
+
+	// Functions used by the parser
+	int addEntry(std::string entryName); // creates a new entry
+	int add(cuobjdumpInst* newInst); //add cuobjdumpInst to list
+	void addConstMemory(int index); // add global const memory
+	void addTex(std::string tex);	// add tex operand
+	bool findEntry(std::string entryName, cuobjdumpEntry& entry); // find and return entry
+
 	void setKernelCount(int k);
 	void readConstMemoryFromElfFile(std::string elf);
-
+	void setLastEntryName(std::string entryName); // sets name of last entry
 	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);
@@ -32,16 +124,22 @@ public:
 
 	void addEntryConstMemory(int index, int entryIndex);
 	void addEntryConstMemory2(char* kernel);
+	void setConstMemoryType(const char* type);
+	void setConstMemoryType2(const char* type); // set type of constant memory
+	void addConstMemoryValue(std::string constMemoryValue); // add const memory
+	void addConstMemoryValue2(std::string constMemoryValue); // add const memory
 	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
-
+	std::list<std::string> getRealTexList(); // get the list of real tex names
+	void setRealTexList(std::list<std::string> realTexList); // set the list of real tex names
+	void printHeaderInstList();
 	void printCuobjdumpLocalMemory();
-
 	void printCuobjdumpInstList();
 	void printCuobjdumpPtxPlusList(cuobjdumpInstList* headerInfo);
 };
 
+#endif //_CUOBJDUMPINSTLIST_H_
diff --git a/cuobjdump_to_ptxplus/cuobjdump_to_ptxplus.cc b/cuobjdump_to_ptxplus/cuobjdump_to_ptxplus.cc
index e8d45ef..fa322e1 100644
--- a/cuobjdump_to_ptxplus/cuobjdump_to_ptxplus.cc
+++ b/cuobjdump_to_ptxplus/cuobjdump_to_ptxplus.cc
@@ -1,7 +1,9 @@
 #include <iostream>
-#include "cuobjdumpInstList.h"
 #include <stdio.h>
 #include <fstream>
+#include <cassert>
+
+#include "cuobjdumpInstList.h"
 
 using namespace std;
 
diff --git a/cuobjdump_to_ptxplus/elf.l b/cuobjdump_to_ptxplus/elf.l
index 3457a7a..f594a00 100644
--- a/cuobjdump_to_ptxplus/elf.l
+++ b/cuobjdump_to_ptxplus/elf.l
@@ -2,7 +2,7 @@
 %{
 #include <stdio.h>
 #include <string.h>
-#include "elf_parser.hpp"
+#include "elf_parser.hh"
 
 #define yylval elf_lval
 
diff --git a/decuda_to_ptxplus/header.l b/cuobjdump_to_ptxplus/header.l
index 5def3bd..170ecf7 100644
--- a/decuda_to_ptxplus/header.l
+++ b/cuobjdump_to_ptxplus/header.l
@@ -32,7 +32,7 @@
 %option prefix="header_"
 %{
 #include <iostream>
-#include "headerParser.hpp"
+#include "header_parser.hh"
 #include <string.h>
 
 /* not sure if this is needed*/
diff --git a/decuda_to_ptxplus/header.y b/cuobjdump_to_ptxplus/header.y
index 6aa545e..1bf62ef 100644
--- a/decuda_to_ptxplus/header.y
+++ b/cuobjdump_to_ptxplus/header.y
@@ -28,12 +28,12 @@
 
 %{
 #include <iostream>
-#include "decudaInstList.h"
+#include "cuobjdumpInstList.h"
 
 int yylex(void);
 void yyerror(const char*);
 
-extern decudaInstList *g_headerList;
+extern cuobjdumpInstList *g_headerList;
 extern void output(const char * text);
 %}
 
@@ -82,9 +82,9 @@ statement	: compilerDirective literal literal	{}
 		| compilerDirective identifierList LEFTPAREN parameterList RIGHTPAREN	{}
 		;
 
-compilerDirective	: DOTVERSION	{ output(".version"); decudaInst *instEntry = new decudaInst(); instEntry->setBase(".version"); g_headerList->add(instEntry);}
-			| DOTTARGET	{ output(".target"); decudaInst *instEntry = new decudaInst(); instEntry->setBase(".target"); g_headerList->add(instEntry);}
-			| DOTENTRY	{ output(".entry"); decudaInst *instEntry = new decudaInst(); instEntry->setBase(".entry"); g_headerList->add(instEntry);}
+compilerDirective	: DOTVERSION	{ output(".version"); cuobjdumpInst *instEntry = new cuobjdumpInst(); instEntry->setBase(".version"); g_headerList->add(instEntry);}
+			| DOTTARGET	{ output(".target"); cuobjdumpInst *instEntry = new cuobjdumpInst(); instEntry->setBase(".target"); g_headerList->add(instEntry);}
+			| DOTENTRY	{ output(".entry"); cuobjdumpInst *instEntry = new cuobjdumpInst(); instEntry->setBase(".entry"); g_headerList->add(instEntry);}
 			;
 
 identifierList	: identifierList IDENTIFER	{ output(" "); output($2); g_headerList->getListEnd().addOperand($2); }
@@ -98,7 +98,7 @@ parameterList	: parameterList parameter
 parameter	: stateSpace opTypes IDENTIFER	{ output(" "); output($3); g_headerList->getListEnd().addOperand($3);}
 		;
 
-stateSpace	: DOTPARAM	{ output("\n.param"); decudaInst *instEntry = new decudaInst(); instEntry->setBase(".param"); g_headerList->add(instEntry); }
+stateSpace	: DOTPARAM	{ output("\n.param"); cuobjdumpInst *instEntry = new cuobjdumpInst(); instEntry->setBase(".param"); g_headerList->add(instEntry); }
 		;
 
 opTypes		: DOTU64		{ output(".u64"); g_headerList->getListEnd().addTypeModifier(".u64");}
diff --git a/cuobjdump_to_ptxplus/ptx_parser.h b/cuobjdump_to_ptxplus/ptx_parser.h
index 5111743..e32f6d8 100644
--- a/cuobjdump_to_ptxplus/ptx_parser.h
+++ b/cuobjdump_to_ptxplus/ptx_parser.h
@@ -7,8 +7,9 @@
 #include <assert.h>
 #include <iostream>
 #include <sstream>
-#include "decudaInstList.h"
 #include <string.h>
+#include <stdio.h>
+#include "cuobjdumpInstList.h"
 
 #define NON_ARRAY_IDENTIFIER 1
 #define ARRAY_IDENTIFIER_NO_DIM 2
@@ -79,9 +80,7 @@ void add_constptr(const char* identifier1, const char* identifier2, int offset)
 
 /*non-dummy stuff below this point*/
 
-
-
-extern decudaInstList *g_headerList;
+extern cuobjdumpInstList *g_headerList;
 
 // Global variable to track if we are currently inside a entry directive
 bool inEntryDirective = false;
@@ -118,13 +117,13 @@ void add_function_name( const char *headerInput )
 void add_space_spec( enum _memory_space_t spec, int value )
 {
 	DPRINTF("spec=%u", spec);
-	decudaInst *instEntry;
+	cuobjdumpInst *instEntry;
 	static int constmemindex=1;
 	switch(spec)
 	{
 		case param_space_unclassified:
 			if(inEntryDirective && inParamDirective) {
-				instEntry = new decudaInst();
+				instEntry = new cuobjdumpInst();
 				instEntry->setBase(".param");
 				g_headerList->add(instEntry);
 			}
@@ -132,7 +131,7 @@ void add_space_spec( enum _memory_space_t spec, int value )
 		case tex_space:
 			inTexDirective = true;
 			/*
-			instEntry = new decudaInst();
+			instEntry = new cuobjdumpInst();
 			instEntry->setBase(".tex");
 			g_headerList->add(instEntry);
 			*/
@@ -141,7 +140,7 @@ void add_space_spec( enum _memory_space_t spec, int value )
 			if(!inEntryDirective) {
 				/*
 				inConstDirective = true;
-				instEntry = new decudaInst();
+				instEntry = new cuobjdumpInst();
 				instEntry->setBase(".const");
 				g_headerList->add(instEntry);
 				*/
@@ -219,7 +218,7 @@ void add_scalar_type_spec( int headerInput )
 void add_version_info( float versionNumber, unsigned ext)
 {
 	DPRINTF("");
-	decudaInst *instEntry = new decudaInst();
+	cuobjdumpInst *instEntry = new cuobjdumpInst();
 	instEntry->setBase(".version");
 	g_headerList->add(instEntry);
 
@@ -236,7 +235,7 @@ void add_version_info( float versionNumber, unsigned ext)
 void target_header(char* firstTarget)
 {
 	DPRINTF("%s", firstTarget);
-	decudaInst *instEntry = new decudaInst();
+	cuobjdumpInst *instEntry = new cuobjdumpInst();
 	instEntry->setBase(".target");
 	g_headerList->add(instEntry);	
 
@@ -246,7 +245,7 @@ void target_header(char* firstTarget)
 void target_header2(char* firstTarget, char* secondTarget)
 {
 	DPRINTF("%s, %s", firstTarget, secondTarget);
-	decudaInst *instEntry = new decudaInst();
+	cuobjdumpInst *instEntry = new cuobjdumpInst();
 	instEntry->setBase(".target");
 	g_headerList->add(instEntry);	
 
@@ -258,7 +257,7 @@ void target_header2(char* firstTarget, char* secondTarget)
 void target_header3(char* firstTarget, char* secondTarget, char* thirdTarget)
 {
 	DPRINTF("%s, %s, %s", firstTarget, secondTarget, thirdTarget);
-	decudaInst *instEntry = new decudaInst();
+	cuobjdumpInst *instEntry = new cuobjdumpInst();
 	instEntry->setBase(".target");
 	g_headerList->add(instEntry);
 
@@ -290,7 +289,7 @@ void func_header(const char* headerBase)
 	if((strcmp(headerBase, ".entry")==0)||(strcmp(headerBase, ".func")==0)) {
 		inEntryDirective = true;
 		g_headerList->addEntry("");
-		decudaInst *instEntry = new decudaInst();
+		cuobjdumpInst *instEntry = new cuobjdumpInst();
 		instEntry->setBase(headerBase);
 		g_headerList->add(instEntry);
 
diff --git a/cuobjdump_to_ptxplus/sass.l b/cuobjdump_to_ptxplus/sass.l
index e665770..859b672 100644
--- a/cuobjdump_to_ptxplus/sass.l
+++ b/cuobjdump_to_ptxplus/sass.l
@@ -1,6 +1,6 @@
 %{
 #include <stdio.h>
-#include "sass_parser.hpp"
+#include "sass_parser.hh"
 #include <string.h>
 
 
diff --git a/decuda_to_ptxplus/stringList.cc b/cuobjdump_to_ptxplus/stringList.cc
index 099e2b9..8d1ae42 100644
--- a/decuda_to_ptxplus/stringList.cc
+++ b/cuobjdump_to_ptxplus/stringList.cc
@@ -115,7 +115,7 @@ bool stringList::remove(int index)
 	return true;
 }
 
-//print out all the Decuda Instructions in the list
+//print out all the cuobjdump Instructions in the list
 void stringList::printStringList()
 {
 	stringListPiece* currentPiece = m_listStart;
diff --git a/decuda_to_ptxplus/stringList.h b/cuobjdump_to_ptxplus/stringList.h
index e1dfaf0..fe0139e 100644
--- a/decuda_to_ptxplus/stringList.h
+++ b/cuobjdump_to_ptxplus/stringList.h
@@ -25,7 +25,8 @@
 // OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
 // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
-
+#ifndef _STRINGLIST_H_
+#define _STRINGLIST_H_
 struct stringListPiece
 {
 	const char* stringText;
@@ -60,3 +61,4 @@ public:
 	//print representation
 	void printStringList();
 };
+#endif //_STRINGLIST_H_
diff --git a/decuda_to_ptxplus/decuda.l b/decuda_to_ptxplus/decuda.l
deleted file mode 100644
index 2c7d2d2..0000000
--- a/decuda_to_ptxplus/decuda.l
+++ /dev/null
@@ -1,301 +0,0 @@
-/*
-// Copyright (c) 2009-2011, Jimmy Kwa,
-// The University of British Columbia
-// All rights reserved.
-//
-// Redistribution and use in source and binary forms, with or without
-// modification, are permitted provided that the following conditions are met:
-//
-// Redistributions of source code must retain the above copyright notice, this
-// list of conditions and the following disclaimer.
-// Redistributions in binary form must reproduce the above copyright notice, this
-// list of conditions and the following disclaimer in the documentation and/or
-// other materials provided with the distribution.
-// Neither the name of The University of British Columbia nor the names of its
-// contributors may be used to endorse or promote products derived from this
-// software without specific prior written permission.
-//
-// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
-// ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
-// WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
-// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
-// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
-// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
-// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
-// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
-// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
-// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-*/
-
-%{
-#include <iostream>
-#include "decudaParser.hpp"
-#include <string.h>
-
-/* not sure if this is needed*/
-#define	YY_NEVER_INTERACTIVE	1
-
-void yyerror(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*/
-\.entry	yylval.string_value = strdup(yytext); return DOTENTRY;
-\.lmem	yylval.string_value = strdup(yytext); return DOTLMEM;
-\.smem	yylval.string_value = strdup(yytext); return DOTSMEM;
-\.reg	yylval.string_value = strdup(yytext); return DOTREG;
-\.bar	yylval.string_value = strdup(yytext); return DOTBAR;
-
-	/*constant memory directives*/
-\.constseg	return DOTCONSTSEG;
-const		return CONST;
-d\.u32		return DDOTU32;
-d\.f32		return DDOTF32;
-
-	/*guard predicate*/
-@!?\$p{numeric}+	yylval.string_value = strdup(yytext); return GUARDPRED;
-
-	/*label*/
-label{numeric}+		yylval.string_value = strdup(yytext); return LABEL;
-
-	/*base instructions*/
-cvt	yylval.string_value = strdup(yytext); return CVT;
-shl	yylval.string_value = strdup(yytext); return SHL;
-shr	yylval.string_value = strdup(yytext); return SHR;
-add	yylval.string_value = strdup(yytext); return ADD;
-addc	yylval.string_value = strdup(yytext); return ADDC;
-sub	yylval.string_value = strdup(yytext); return SUB;
-subr	yylval.string_value = strdup(yytext); return SUBR;
-movsh	yylval.string_value = strdup(yytext); return MOVSH;
-mov	yylval.string_value = strdup(yytext); return MOV;
-bar\.sync	yylval.string_value = strdup(yytext); return BARSYNC;
-mul	yylval.string_value = strdup(yytext); return MUL;
-mad24	yylval.string_value = strdup(yytext); return MAD24;
-mad24c1	yylval.string_value = strdup(yytext); return MAD24C1;
-set	yylval.string_value = strdup(yytext); return SET;
-mad	yylval.string_value = strdup(yytext); return MAD;
-mul24	yylval.string_value = strdup(yytext); return MUL24;
-set\?{numeric}+\?	yylval.string_value = strdup(yytext); return SETUNKNOWN;
-rcp	yylval.string_value = strdup(yytext); return RCP;
-pre\.ex2	yylval.string_value = strdup(yytext); return PREEX2;
-pre\.sin	yylval.string_value = strdup(yytext); return PRESIN;
-ex2	yylval.string_value = strdup(yytext); return EX2;
-sin	yylval.string_value = strdup(yytext); return SIN;
-cos	yylval.string_value = strdup(yytext); return COS;
-lg2	yylval.string_value = strdup(yytext); return LG2;
-rsqrt	yylval.string_value = strdup(yytext); return RSQRT;
-mac	yylval.string_value = strdup(yytext); return MAC;
-and	yylval.string_value = strdup(yytext); return AND;
-andn	yylval.string_value = strdup(yytext); return ANDN;
-nandn	yylval.string_value = strdup(yytext); return NANDN;
-xor	yylval.string_value = strdup(yytext); return XOR;
-or	yylval.string_value = strdup(yytext); return OR;
-orn	yylval.string_value = strdup(yytext); return ORN;
-min	yylval.string_value = strdup(yytext); return MIN;
-tex	yylval.string_value = strdup(yytext); return TEX;
-bra\.label	yylval.string_value = strdup(yytext); return BRA;
-join\.label	yylval.string_value = strdup(yytext); return JOIN;
-breakaddr.label	yylval.string_value = strdup(yytext); return BREAKADDR;
-nop\.end	yylval.string_value = strdup(yytext); return NOPEND;
-nop\.join	yylval.string_value = strdup(yytext); return NOPJOIN;
-nop	yylval.string_value = strdup(yytext); return NOP;
-call.label	yylval.string_value = strdup(yytext); return CALL;
-not	yylval.string_value = strdup(yytext); return NOT;
-max	yylval.string_value = strdup(yytext); return MAX;
-delta	yylval.string_value = strdup(yytext); return DELTA;
-break	yylval.string_value = strdup(yytext); return BREAK;
-inc	yylval.string_value = strdup(yytext); return INC;
-exch	yylval.string_value = strdup(yytext); return EXCH;
-cas	yylval.string_value = strdup(yytext); return CAS;
-norn	yylval.string_value = strdup(yytext); return NORN;
-nxor	yylval.string_value = strdup(yytext); return NXOR;
-sad		yylval.string_value = strdup(yytext); return SAD;
-op\.13	yylval.string_value = strdup(yytext); return OP13;
-op\.e5	yylval.string_value = strdup(yytext); return OPE5;
-op\.e6	yylval.string_value = strdup(yytext); return OPE6;
-op\.d0	yylval.string_value = strdup(yytext); return OPD0;
-
-return	yylval.string_value = strdup(yytext); return RETURN;
-
-	/*instruction modifiers*/
-\.rp	yylval.string_value = strdup(yytext); return DOTRP;
-\.rm	yylval.string_value = strdup(yytext); return DOTRM;
-\.rn	yylval.string_value = strdup(yytext); return DOTRN;
-\.rni	yylval.string_value = strdup(yytext); return DOTRNI;
-\.rz	yylval.string_value = strdup(yytext); return DOTRZ;
-\.rzi	yylval.string_value = strdup(yytext); return DOTRZI;
-\.half	yylval.string_value = strdup(yytext); return DOTHALF;
-\.end	yylval.string_value = strdup(yytext); return DOTEND;
-\.lo	yylval.string_value = strdup(yytext); return DOTLO;
-\.abs	yylval.string_value = strdup(yytext); return DOTABS;
-\.join	yylval.string_value = strdup(yytext); return DOTJOIN;
-\.neg	yylval.string_value = strdup(yytext); return DOTNEG;
-\.sat	yylval.string_value = strdup(yytext); return DOTSAT;
-
-	/* comparison operators */
-\.eq	yylval.string_value = strdup(yytext); return DOTEQ;
-\.ne	yylval.string_value = strdup(yytext); return DOTNE;
-\.lt	yylval.string_value = strdup(yytext); return DOTLT;
-\.le	yylval.string_value = strdup(yytext); return DOTLE;
-\.gt	yylval.string_value = strdup(yytext); return DOTGT;
-\.ge	yylval.string_value = strdup(yytext); return DOTGE;
-\.ls	yylval.string_value = strdup(yytext); return DOTLS;
-\.hi	yylval.string_value = strdup(yytext); return DOTHI;
-\.hs	yylval.string_value = strdup(yytext); return DOTHS;
-\.equ	yylval.string_value = strdup(yytext); return DOTEQU;
-\.neu	yylval.string_value = strdup(yytext); return DOTNEU;
-\.ltu	yylval.string_value = strdup(yytext); return DOTLTU;
-\.leu	yylval.string_value = strdup(yytext); return DOTLEU;
-\.gtu	yylval.string_value = strdup(yytext); return DOTGTU;
-\.geu	yylval.string_value = strdup(yytext); return DOTGEU;
-\.num	yylval.string_value = strdup(yytext); return DOTNUM;
-\.nan	yylval.string_value = strdup(yytext); return DOTNAN;
-\.nsf	yylval.string_value = strdup(yytext); return DOTNSF;
-\.sf	yylval.string_value = strdup(yytext); return DOTSF;
-\.cf	yylval.string_value = strdup(yytext); return DOTCF;
-
-	/*operand types*/
-\.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;
-\.u8	yylval.string_value = strdup(yytext); return DOTU8;
-\.u16	yylval.string_value = strdup(yytext); return DOTU16;
-\.u32	yylval.string_value = strdup(yytext); return DOTU32;
-\.u64	yylval.string_value = strdup(yytext); return DOTU64;
-\.b8	yylval.string_value = strdup(yytext); return DOTB8;
-\.b16	yylval.string_value = strdup(yytext); return DOTB16;
-\.b32	yylval.string_value = strdup(yytext); return DOTB32;
-\.b64	yylval.string_value = strdup(yytext); return DOTB64;
-\.b128	yylval.string_value = strdup(yytext); return DOTB128;
-\.f16	yylval.string_value = strdup(yytext); return DOTF16;
-\.f32	yylval.string_value = strdup(yytext); return DOTF32;
-\.f64	yylval.string_value = strdup(yytext); return DOTF64;
-
-	/*geometries*/
-\.1d	yylval.string_value = strdup(yytext); return DOT1D;
-\.2d	yylval.string_value = strdup(yytext); return DOT2D;
-\.3d	yylval.string_value = strdup(yytext); return DOT3D;
-
-	/*register operands*/
-[\-]?\$r{numeric}+	yylval.string_value = strdup(yytext); return REGISTER;
-[\-]?\$r{numeric}+\.lo	yylval.string_value = strdup(yytext); return REGISTERLO;
-[\-]?\$r{numeric}+\.hi	yylval.string_value = strdup(yytext); return REGISTERHI;
-[\-]?\$ofs{numeric}+	yylval.string_value = strdup(yytext); return OFFSETREGISTER;
-\$p{numeric}+		yylval.string_value = strdup(yytext); return PREDREGISTER;
-\$o{numeric}+		yylval.string_value = strdup(yytext); return PREDREGISTER2;
-
-	/*tex operand*/
-[\-]?\$tex{numeric}+	yylval.string_value = strdup(yytext); return TEXOP;
-
-	/*vector operands*/
-\{($r{numeric}+|_)\}	yylval.string_value = strdup(yytext); return VECTOR1D;
-\{($r{numeric}+|_)\,($r{numeric}+|_)\}	yylval.string_value = strdup(yytext); return VECTOR2D;
-\{($r{numeric}+|_)\,($r{numeric}+|_)\,($r{numeric}+|_)\}	yylval.string_value = strdup(yytext); return VECTOR3D;
-\{($r{numeric}+|_)\,($r{numeric}+|_)\,($r{numeric}+|_)\,($r{numeric}+|_)\}	yylval.string_value = strdup(yytext); return VECTOR4D;
-
-	/*special registers*/
-"%ntid.x"	return NTIDXREGISTER;
-"%ntid.y"	return NTIDYREGISTER;
-"%ctaid.x"	return CTAIDXREGISTER;
-"%nctaid.x"	return NCTAIDXREGISTER;
-"%ctaid.y"	return CTAIDYREGISTER;
-"%nctaid.y"	return NCTAIDYREGISTER;
-"%clock"	return CLOCKREGISTER;
-
-	/*memory location operands*/
-[\-]?l\[{hexliteral}\]				yylval.string_value = strdup(yytext); return LMEMLOCATION;
-[\-]?l\[\$r{numeric}+\]				yylval.string_value = strdup(yytext); return LMEMLOCATION;
-[\-]?l\[\$ofs{numeric}+\+{hexliteral}\]		yylval.string_value = strdup(yytext); return LMEMLOCATION;
-[\-]?l\[\$ofs{numeric}+\+={hexliteral}\]	yylval.string_value = strdup(yytext); return LMEMLOCATION;
-
-[\-]?s\[{hexliteral}\]				yylval.string_value = strdup(yytext); return SMEMLOCATION;
-[\-]?s\[\$r{numeric}+\]				yylval.string_value = strdup(yytext); return SMEMLOCATION;
-[\-]?s\[\$ofs{numeric}+\+{hexliteral}\]		yylval.string_value = strdup(yytext); return SMEMLOCATION;
-[\-]?s\[\$ofs{numeric}+\+={hexliteral}\]	yylval.string_value = strdup(yytext); return SMEMLOCATION;
-
-[\-]?g\[{hexliteral}\]				yylval.string_value = strdup(yytext); return GMEMLOCATION;
-[\-]?g\[\$r{numeric}+\]				yylval.string_value = strdup(yytext); return GMEMLOCATION;
-[\-]?g\[\$ofs{numeric}+\+{hexliteral}\]		yylval.string_value = strdup(yytext); return GMEMLOCATION;
-[\-]?g\[\$ofs{numeric}+\+={hexliteral}\]	yylval.string_value = strdup(yytext); return GMEMLOCATION;
-
-
-[\-]?c{numeric}+\[{hexliteral}\]			yylval.string_value = strdup(yytext); return CMEMLOCATION;
-[\-]?c{numeric}+\[\$r{numeric}+\]			yylval.string_value = strdup(yytext); return CMEMLOCATION;
-[\-]?c{numeric}+\[\$ofs{numeric}+\+{hexliteral}\]	yylval.string_value = strdup(yytext); return CMEMLOCATION;
-[\-]?c{numeric}+\[\$ofs{numeric}+\+={hexliteral}\]	yylval.string_value = strdup(yytext); return CMEMLOCATION;
-
-
-	/*identifier + literals*/
-{identifier}	yylval.string_value = strdup(yytext); return IDENTIFER;
-{hexliteral}	yylval.string_value = strdup(yytext); return HEXLITERAL; //change to int later?
-{octliteral}	yylval.string_value = strdup(yytext); return OCTLITERAL; //change to int later?
-{binliteral}	yylval.string_value = strdup(yytext); return BINLITERAL; //change to int later?
-{decliteral}	yylval.string_value = strdup(yytext); return DECLITERAL; //change to int later?
-{fltliteral}	yylval.string_value = strdup(yytext); return FLTLITERAL;
-
-\:	return COLON;
-
-{pipe}	return PIPE;
-
-	/*braces*/
-\{	yylval.string_value = strdup(yytext); return LEFTBRACE;
-\}	yylval.string_value = strdup(yytext); return RIGHTBRACE;
-
-	/* pound sign*/
-"#"	yylval.string_value = strdup(yytext); return POUND;
-
-	/* comma*/
-","	yylval.string_value = strdup(yytext); return COMMA;
-
-	/* underscore sign*/
-"_"	yylval.string_value = strdup(yytext); return UNDERSCORE;
-
-	/*Need to create to newline tokens*/
-{newline}+	return NEWLINE;
-
-
-	/*discard whitespace*/
-{whitespace}+
-	/*discard comments*/
-"//"{anycharbutnewline}*
-
-
-	/*reached end of file*/
-<<EOF>>	return 0;
-
-	/*None of the rules matched. Throw error*/
-.			yyerror("Invalid token");
-
-%%
-
-void yyerror(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/decuda_to_ptxplus/decuda.y b/decuda_to_ptxplus/decuda.y
deleted file mode 100644
index 0b297f8..0000000
--- a/decuda_to_ptxplus/decuda.y
+++ /dev/null
@@ -1,339 +0,0 @@
-// Copyright (c) 2009-2011, Jimmy Kwa,
-// The University of British Columbia
-// All rights reserved.
-//
-// Redistribution and use in source and binary forms, with or without
-// modification, are permitted provided that the following conditions are met:
-//
-// Redistributions of source code must retain the above copyright notice, this
-// list of conditions and the following disclaimer.
-// Redistributions in binary form must reproduce the above copyright notice, this
-// list of conditions and the following disclaimer in the documentation and/or
-// other materials provided with the distribution.
-// Neither the name of The University of British Columbia nor the names of its
-// contributors may be used to endorse or promote products derived from this
-// software without specific prior written permission.
-//
-// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
-// ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
-// WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
-// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
-// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
-// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
-// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
-// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
-// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
-// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-
-%{
-#include <iostream>
-#include "decudaInstList.h"
-
-int yylex(void);
-void yyerror(const char*);
-
-extern decudaInstList *g_instList;
-
-decudaInst *instEntry;
-%}
-
-
-%union {
-  double double_value;
-  float  float_value;
-  int    int_value;
-  char * string_value;
-  void * ptr_value;
-}
-
-%token <string_value> DOTENTRY DOTLMEM DOTSMEM DOTREG DOTBAR
-
-%token <string_value> GUARDPRED LABEL
-
-%token <string_value> CVT SHL SHR ADD SUB SUBR MOVSH MOV BARSYNC MUL MAD24 MAD24C1 SET RETURN MAD MUL24 SETUNKNOWN RCP PREEX2 PRESIN EX2 SIN COS LG2 RSQRT MAC BRA NOPEND JOIN NOPJOIN NOP AND ANDN XOR MIN BREAKADDR TEX CALL NOT OR MAX DELTA BREAK INC EXCH CAS NORN ADDC ORN NANDN NXOR SAD OPE5 OPE6 OP13 OPD0
-
-%token <string_value> DOTEQ DOTNE DOTLT DOTLE DOTGT DOTGE DOTLO DOTLS DOTHI DOTHS DOTEQU DOTNEU DOTLTU DOTLEU DOTGTU DOTGEU DOTNUM DOTNAN DOTNSF DOTSF DOTCF
-
-%token <string_value> DOTRP DOTRM DOTRN DOTRNI DOTRZ DOTRZI DOTHALF DOTEND DOTABS DOTJOIN DOTNEG DOTSAT
-
-%token <string_value> DOT1D DOT2D DOT3D
-
-%token <string_value> DOTS8 DOTS16 DOTS32 DOTS64
-%token <string_value> DOTU8 DOTU16 DOTU32 DOTU64
-%token <string_value> DOTB8 DOTB16 DOTB32 DOTB64 DOTB128
-%token <string_value> DOTF16 DOTF32 DOTF64
-
-%token <string_value> REGISTER REGISTERLO REGISTERHI OFFSETREGISTER PREDREGISTER PREDREGISTER2
-%token <string_value> VECTOR1D VECTOR2D VECTOR3D VECTOR4D
-
-%token <string_value> TEXOP
-
-%token <string_value> NTIDXREGISTER NTIDYREGISTER CTAIDXREGISTER NCTAIDXREGISTER CTAIDYREGISTER NCTAIDYREGISTER CLOCKREGISTER
-
-%token <string_value> LMEMLOCATION SMEMLOCATION GMEMLOCATION CMEMLOCATION
-
-%token <string_value> IDENTIFER
-
-%token <string_value> DOTCONSTSEG CONST DDOTU32 DDOTF32;
-
-	/*change these 4 to int later?*/
-%token <string_value> HEXLITERAL
-%token <string_value> OCTLITERAL
-%token <string_value> BINLITERAL
-%token <string_value> DECLITERAL
-%token <string_value> FLTLITERAL
-
-%token <string_value> COLON
-%token <string_value> LEFTBRACE
-%token <string_value> RIGHTBRACE
-%token NEWLINE
-%token PIPE
-%token <string_value> COMMA
-%token <string_value> POUND
-%token <string_value> UNDERSCORE
-
-
-	/* set types for rules */
-%type<string_value> registerlocation
-%type<string_value> predicate predicateModifier
-%type<string_value> simpleInstructions
-%type<string_value> comparisonOp
-%type<string_value> constMemoryTypes
-%type<string_value> compilerDirective
-
-%%
-
-	/*translation rules*/
-program		: statementList			{ printf("No parsing errors\n");  }
-		;
-
-statementList	: statementList statement newlines	{ printf("\n"); }
-		| statement newlines			{ printf("\n"); }
-		| newlines				{}
-		;
-
-newlines	: NEWLINE
-		| newlines NEWLINE;
-
-statement	: LEFTBRACE			{ printf("{"); }
-		| RIGHTBRACE			{ printf("}"); }
-		| compilerDirectiveStatement	{}
-		| constMemoryDirectiveStatement	{}
-		| { instEntry = new decudaInst(); }	assemblyInstruction
-		;
-
-compilerDirectiveStatement	: DOTENTRY IDENTIFER		{ printf("%s %s", $1, $2);
-								  g_instList->addEntry($2);
-								  instEntry = new decudaInst(); instEntry->setBase($1); 					  			  g_instList->add(instEntry);
-								  g_instList->getListEnd().addOperand($2);}
-
-				| compilerDirective 		{ }
-				;
-
-compilerDirective	: DOTLMEM DECLITERAL	{ printf(".lmem "); printf($2); g_instList->setLastEntryLMemSize(atoi($2));}
-			| DOTSMEM DECLITERAL	{ printf(".smem "); printf($2); }
-			| DOTREG DECLITERAL	{ printf(".reg "); printf($2); }
-			| DOTBAR DECLITERAL	{ printf(".bar "); printf($2); }
-			;
-
-assemblyInstruction	: baseInstruction modifierList operandList	{ }
-			| otherInstruction	{ }
-			| predicate assemblyInstruction	{ }
-			| instructionLabel assemblyInstruction	{ }
-			;
-
-instructionLabel	: LABEL COLON	{ printf($1); printf(": "); instEntry->setLabel($1); };
-
-predicate		: GUARDPRED	{ printf($1); printf(" "); instEntry->setPredicate($1); }
-			| GUARDPRED predicateModifier	{ printf($1); printf($2); printf(" "); instEntry->setPredicate($1); instEntry->addPredicateModifier($2);}
-			;
-
-predicateModifier	: DOTNE 	{ }
-			| DOTNEU	{ }
-			| DOTEQ		{ }
-			| DOTEQU	{ }
-			| DOTGTU	{ }
-			| DOTLTU	{ }
-			| DOTLE		{ }
-			| DOTGE		{ }
-			| DOTNSF	{ }
-			| DOTSF		{ }
-			| DOTCF		{ }
-			;
-
-baseInstruction		: simpleInstructions	{ printf($1); instEntry->setBase($1); g_instList->add(instEntry);}
-			| SET comparisonOp { printf($1); instEntry->setBase($1); g_instList->add(instEntry); printf($2); instEntry->addBaseModifier($2); }
-			| SET DOTJOIN comparisonOp { printf($1); instEntry->setBase($1); g_instList->add(instEntry); printf($2); instEntry->addBaseModifier($2); instEntry->addBaseModifier($3); }
-			;
-
-simpleInstructions	: CVT | SHL | SHR | ADD | SUB | SUBR | MOVSH | MOV | BARSYNC | MUL | MAD24 | MAD24C1 | MAD | MUL24 | SETUNKNOWN | RCP | PREEX2 | PRESIN | SIN | COS | EX2 | LG2 | RSQRT | MAC | AND | ANDN | XOR | MIN | TEX | NOT | OR | MAX | DELTA | INC | EXCH | CAS | NORN | ADDC | ORN | NANDN | NXOR | SAD | OPD0
-			;
-
-otherInstruction	: RETURN	{ printf($1); instEntry->setBase($1); g_instList->add(instEntry);}
-			| BRA LABEL 	{ printf($1); printf(" "); printf($2); instEntry->setBase($1); g_instList->add(instEntry); instEntry->addOperand($2);  }
-			| JOIN LABEL 	{ printf($1); printf(" "); printf($2); instEntry->setBase($1); g_instList->add(instEntry); instEntry->addOperand($2);  }
-			| BREAKADDR LABEL 	{ printf($1); printf(" "); printf($2); instEntry->setBase($1); g_instList->add(instEntry); instEntry->addOperand($2);  }
-			| CALL LABEL 	{ printf($1); printf(" "); printf($2); instEntry->setBase($1); g_instList->add(instEntry); instEntry->addOperand($2);  }
-			| NOPEND	{ printf($1); instEntry->setBase($1); g_instList->add(instEntry); }
-			| NOPJOIN	{ printf($1); instEntry->setBase($1); g_instList->add(instEntry); }
-			| NOP		{ printf($1); instEntry->setBase($1); g_instList->add(instEntry); }
-			| BREAK		{ printf($1); instEntry->setBase($1); g_instList->add(instEntry); }
-			| OP13 PREDREGISTER		{ printf($1); instEntry->setBase($1); g_instList->add(instEntry); g_instList->addPredicate($2);}
-			| OPE6 { printf($1); instEntry->setBase($1); g_instList->add(instEntry); }
-			| OPE5 { printf($1); instEntry->setBase($1); g_instList->add(instEntry); }
-
-comparisonOp	: DOTEQ			{ }
-		| DOTNE			{ }
-		| DOTLE			{ }
-		| DOTLT			{ }
-		| DOTGE			{ }
-		| DOTGT			{ }
-		| DOTLO			{ }
-		| DOTLS			{ }
-		| DOTHI			{ }
-		| DOTHS			{ }
-		| DOTEQU		{ }
-		| DOTNEU		{ }
-		| DOTLTU		{ }
-		| DOTLEU		{ }
-		| DOTGTU		{ }
-		| DOTGEU		{ }
-		| DOTNUM		{ }
-		| DOTNAN		{ }
-		;
-
-modifierList	: modifierList modifier	{}
-		| modifier		{}
-		;
-
-modifier	: opTypes		{}
-		| geometries		{}
-		| DOTRP			{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
-		| DOTRM			{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
-		| DOTRN			{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
-		| DOTRNI		{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
-		| DOTRZ			{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
-		| DOTRZI		{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
-		| DOTHALF		{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
-		| DOTEND		{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
-		| DOTLO			{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
-		| DOTABS		{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
-		| DOTJOIN		{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
-		| DOTNEG		{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
-		| DOTSAT		{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
-		;
-
-
-geometries	: DOT1D			{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
-		| DOT2D			{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
-		| DOT3D			{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
-		;
-
-opTypes		: 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);}
-		| DOTU8			{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
-		| DOTU16		{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
-		| DOTU32		{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
-		| DOTU64		{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
-		| DOTB8			{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
-		| DOTB16		{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
-		| DOTB32		{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
-		| DOTB64		{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
-		| DOTB128		{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
-		| DOTF16		{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
-		| DOTF32		{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
-		| DOTF64		{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
-		;
-
-operandList	: operandList COMMA operand	{}
-		| operand
-		;
-
-operand		: { printf(" "); } registerlocation
-		| { printf(" "); } specialRegister
-		| { printf(" "); } memorylocation
-		| { printf(" "); } immediateValue	
-		;
-
-registerlocation	: REGISTER		{ printf($1); g_instList->addRegister($1);}
-			| vector		{ }
-			| REGISTERLO		{ printf($1); g_instList->addRegister($1,true);}
-			| REGISTERHI		{ printf($1); g_instList->addRegister($1,true);}
-			| OFFSETREGISTER	{ printf($1); g_instList->addRegister($1);}
-			| PREDREGISTER 		{ printf($1); g_instList->addPredicate($1);}
-			| PREDREGISTER PIPE PREDREGISTER2
-						{ printf($1); printf("|"); printf($3); g_instList->addDoublePredReg($1,$3);}
-			| PREDREGISTER PIPE REGISTER
-						{ printf($1); printf("|"); printf($3); g_instList->addDoublePredReg($1,$3);}
-			| TEXOP			{ printf($1); g_instList->addTex($1);}
-			;
-
-vector		: VECTOR1D	{ printf($1); g_instList->addVector($1,1); }
-		| VECTOR2D	{ printf($1); g_instList->addVector($1,2); }
-		| VECTOR3D	{ printf($1); g_instList->addVector($1,3); }
-		| VECTOR4D	{ printf($1); g_instList->addVector($1,4); }
-		;
-
-specialRegister	: NTIDXREGISTER		{ printf("%%ntid.x"); g_instList->getListEnd().addOperand("%%ntid.x");}
-		| NTIDYREGISTER		{ printf("%%ntid.y"); g_instList->getListEnd().addOperand("%%ntid.y");}
-		| CTAIDXREGISTER	{ printf("%%ctaid.x"); g_instList->getListEnd().addOperand("%%ctaid.x");}
-		| CTAIDYREGISTER	{ printf("%%ctaid.y"); g_instList->getListEnd().addOperand("%%ctaid.y");}
-		| NCTAIDXREGISTER	{ printf("%%nctaid.x"); g_instList->getListEnd().addOperand("%%nctaid.x");}
-		| NCTAIDYREGISTER	{ printf("%%nctaid.y"); g_instList->getListEnd().addOperand("%%nctaid.y");}
-		| CLOCKREGISTER		{ printf("%%clock"); g_instList->getListEnd().addOperand("%%clock");}
-		;
-
-memorylocation	: LMEMLOCATION		{ printf($1); g_instList->addMemoryOperand($1,3);}
-		| SMEMLOCATION		{ printf($1); g_instList->addMemoryOperand($1,1);}
-		| GMEMLOCATION		{ printf($1); g_instList->addMemoryOperand($1,2);}
-		| CMEMLOCATION		{ printf($1); g_instList->addMemoryOperand($1,0);}
-		;
-
-immediateValue	: IDENTIFER	{ printf($1); g_instList->getListEnd().addOperand($1);}
-		| HEXLITERAL	{ printf($1); g_instList->getListEnd().addOperand($1);}
-		| OCTLITERAL	{ printf($1); g_instList->getListEnd().addOperand($1);}
-		| BINLITERAL	{ printf($1); g_instList->getListEnd().addOperand($1);}
-		| DECLITERAL	{ printf($1); g_instList->getListEnd().addOperand($1);}
-		;
-
-
-constMemoryDirectiveStatement 	: POUND DOTCONSTSEG DECLITERAL COLON HEXLITERAL CONST NEWLINE
-				   { g_instList->addEntryConstMemory(atoi($3)); printf("//c"); printf($3); printf(" = "); }
-				  POUND LEFTBRACE NEWLINE
-				  constMemoryStatements
-				  POUND RIGHTBRACE
-				;
-
-constMemoryStatements	: constMemoryStatement
-			| constMemoryStatements constMemoryStatement
-			;
-
-constMemoryStatement	: POUND constMemoryTypes constMemoryList { g_instList->setConstMemoryType($2); } NEWLINE
-
-constMemoryTypes	: DDOTU32 { char* tempString = new char[5];
-					strcpy(tempString, ".u32");
-					$$ = tempString; }
-			| DDOTF32 { char* tempString = new char[5];
-					strcpy(tempString, ".f32");
-					$$ = tempString; }
-			;
-
-constMemoryList		: constMemory
-			| constMemoryList COMMA constMemory
-			;
-
-constMemory		: HEXLITERAL	{printf($1); printf(" "); g_instList->addConstMemoryValue($1); }
-			| OCTLITERAL	{printf($1); printf(" "); g_instList->addConstMemoryValue($1); }
-			| BINLITERAL	{printf($1); printf(" "); g_instList->addConstMemoryValue($1); }
-			| DECLITERAL	{printf($1); printf(" "); g_instList->addConstMemoryValue($1); }
-			| FLTLITERAL	{printf($1); printf(" "); g_instList->addConstMemoryValue($1); }
-			;
-                       
-
-
-%%
-
-/*support c++ functions go here*/
diff --git a/decuda_to_ptxplus/decudaInst.cc b/decuda_to_ptxplus/decudaInst.cc
deleted file mode 100644
index 4f1b9de..0000000
--- a/decuda_to_ptxplus/decudaInst.cc
+++ /dev/null
@@ -1,1523 +0,0 @@
-// Copyright (c) 2009-2011, Jimmy Kwa,
-// The University of British Columbia
-// All rights reserved.
-//
-// Redistribution and use in source and binary forms, with or without
-// modification, are permitted provided that the following conditions are met:
-//
-// Redistributions of source code must retain the above copyright notice, this
-// list of conditions and the following disclaimer.
-// Redistributions in binary form must reproduce the above copyright notice, this
-// list of conditions and the following disclaimer in the documentation and/or
-// other materials provided with the distribution.
-// Neither the name of The University of British Columbia nor the names of its
-// contributors may be used to endorse or promote products derived from this
-// software without specific prior written permission.
-//
-// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
-// ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
-// WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
-// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
-// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
-// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
-// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
-// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
-// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
-// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-
-#include "decudaInst.h"
-#include <iostream>
-#include <string>
-#include <cstring>
-#include <stdlib.h>
-#include <stdio.h>
-
-extern void output(const char * text);
-
-//Constructor
-decudaInst::decudaInst()
-{
-	//initilize everything to empty
-	m_label = "";
-	m_predicate = new stringList();
-	m_base = "";
-	m_baseModifiers = new stringList();
-	m_typeModifiers = new stringList();
-	m_operands = new stringList();
-	m_predicateModifiers = new stringList();
-
-	m_nextDecudaInst = NULL;
-
-   // Set operations per cycle to 8 by default (no penalty)
-   m_opPerCycle = 8;
-}
-
-//retreive instruction mnemonic
-const char* decudaInst::getBase()
-{
-	return m_base;
-}
-
-stringList* decudaInst::getOperands()
-{
-	return m_operands;
-}
-
-stringList* decudaInst::getBaseModifiers()
-{
-	return m_baseModifiers;
-}
-
-stringList* decudaInst::getTypeModifiers()
-{
-        return m_typeModifiers;
-}
-
-//get next instruction in linked list
-//direction is m_listStart to m_listEnd
-decudaInst* decudaInst::getNextDecudaInst()
-{
-	return m_nextDecudaInst;
-}
-
-void decudaInst::setBase(const char* setBaseValue)
-{
-	m_base = setBaseValue;
-}
-
-void decudaInst::addBaseModifier(const char* addBaseMod)
-{
-	stringListPiece* tempPiece = new stringListPiece;
-	tempPiece->stringText = addBaseMod;
-
-	m_baseModifiers->add(tempPiece);
-}
-
-void decudaInst::addTypeModifier(const char* addTypeMod)
-{
-	stringListPiece* tempPiece = new stringListPiece;
-	tempPiece->stringText = addTypeMod;
-
-	m_typeModifiers->add(tempPiece);
-}
-
-void decudaInst::addOperand(const char* addOp)
-{
-	stringListPiece* tempPiece = new stringListPiece;
-	tempPiece->stringText = addOp;
-
-	m_operands->add(tempPiece);
-}
-
-void decudaInst::setPredicate(const char* setPredicateValue)
-{
-	stringListPiece* tempPiece = new stringListPiece;
-	tempPiece->stringText = setPredicateValue;
-
-	m_predicate->add(tempPiece);
-}
-
-void decudaInst::addPredicateModifier(const char* addPredicateMod)
-{
-	stringListPiece* tempPiece = new stringListPiece;
-	tempPiece->stringText = addPredicateMod;
-
-	m_predicateModifiers->add(tempPiece);
-}
-
-void decudaInst::setLabel(const char* setLabelValue)
-{
-	m_label = setLabelValue;
-}
-
-//set next instruction in linked list
-//direction is m_listStart to m_listEnd
-void decudaInst::setNextDecudaInst(decudaInst* setDecudaInstValue)
-{
-	m_nextDecudaInst = setDecudaInstValue;
-}
-
-// returns true if current instruction is start of an entry (i.e. '{')
-bool decudaInst::isEntryStart()
-{
-	return (strcmp(m_base, "{")==0);
-}
-
-//print out .version and .target header lines
-bool decudaInst::printHeaderInst()
-{
-	if(strcmp(m_base, ".version")==0)
-	{
-		output(m_base); output(" ");
-
-		stringListPiece* currentPiece = m_operands->getListStart();
-		output(currentPiece->stringText);
-		currentPiece = currentPiece->nextString;
-
-		if(currentPiece!=NULL)
-		{
-			output("."); output(currentPiece->stringText);
-		}
-      output("+");
-		output("\n");
-	}
-	else if(strcmp(m_base, ".target")==0)
-	{
-		output(m_base); output(" ");
-
-		stringListPiece* currentPiece = m_operands->getListStart();
-		output(currentPiece->stringText);
-		currentPiece = currentPiece->nextString;
-
-		while(currentPiece!=NULL)
-		{
-			output(", "); output(currentPiece->stringText);
-			currentPiece = currentPiece->nextString;
-		}
-		output("\n");
-	}
-	else if(strcmp(m_base, ".tex")==0)
-	{
-		output(m_base); output(" ");
-
-		stringListPiece* currentPiece;
-
-
-		currentPiece = m_baseModifiers->getListStart();
-		output(currentPiece->stringText); output(" ");
-		currentPiece = currentPiece->nextString;
-
-		while(currentPiece!=NULL)
-		{
-			output(" "); output(currentPiece->stringText);
-			currentPiece = currentPiece->nextString;
-		}
-
-		currentPiece = m_operands->getListStart();
-		output(currentPiece->stringText);
-		currentPiece = currentPiece->nextString;
-
-		while(currentPiece!=NULL)
-		{
-			output(" "); output(currentPiece->stringText);
-			currentPiece = currentPiece->nextString;
-		}
-		output(";\n");
-	}
-	else
-	{
-		return false;
-	}
-	return true;
-}
-
-//print out parameters
-bool decudaInst::printHeaderInst2()
-{
-	if(strcmp(m_base, ".param")==0)
-	{
-		output(m_base);
-
-		stringListPiece* currentPiece = m_typeModifiers->getListStart();
-		
-		for(int i=0; (i<m_typeModifiers->getSize())&&(currentPiece!=NULL); i++)
-		{
-			output(" "); output(currentPiece->stringText);
-
-			currentPiece = currentPiece->nextString;
-		}
-
-		currentPiece = m_operands->getListStart();
-		for(int i=0; (i<m_operands->getSize())&&(currentPiece!=NULL); i++)
-		{
-			output(" "); output(currentPiece->stringText);
-
-			currentPiece = currentPiece->nextString;
-		}
-	}
-	else
-	{
-		return false;
-	}
-	return true;
-}
-//print out the Decuda instruction
-void decudaInst::printDecudaInst()
-{
-	/*TODO: print label here*/
-	/*TODO: print predicate here*/
-	/*std::cout << "Instruction Base: " << m_base << "\n";
-
-	std::cout << "Instruction Modifiers: ";
-	m_baseModifiers->printStringList();
-	std::cout << "\n";
-
-	std::cout << "Operand types: ";
-	m_typeModifiers->printStringList();
-	std::cout << "\n";
-
-	std::cout << "Operands: ";
-	m_operands->printStringList();
-	std::cout << "\n\n";*/
-
-	std::cout << m_base << " ";
-	m_baseModifiers->printStringList();
-	std::cout << " ";
-	m_typeModifiers->printStringList();
-	std::cout << " ";
-	m_operands->printStringList();
-	std::cout << "\n";
-
-
-}
-
-
-// Just prints the base and operands
-void decudaInst::printHeaderPtx()
-{
-	output(m_base); output(" ");
-
-	stringListPiece* currentPiece;
-
-	currentPiece = m_baseModifiers->getListStart();
-	for(int i=0; (i<m_baseModifiers->getSize())&&(currentPiece!=NULL); i++)
-	{
-		output(" "); output(currentPiece->stringText);
-		currentPiece = currentPiece->nextString;
-	}
-
-	currentPiece = m_operands->getListStart();
-	for(int i=0; (i<m_operands->getSize())&&(currentPiece!=NULL); i++)
-	{
-		output(" "); output(currentPiece->stringText);
-		currentPiece = currentPiece->nextString;
-	}
-}
-
-
-
-// Print unmodified instruction
-void decudaInst::printDefaultPtx()
-{
-	printLabel();
-	printPredicate();
-
-	output(m_base);
-
-	printBaseModifiers();
-	printTypeModifiers();
-	printOperands();
-
-	output(";");
-}
-
-
-// Print unmodified base modifiers, operands, labels and predicates
-void decudaInst::printLabel()
-{
-	if(m_label != "") {
-		output(m_label);
-		output(": ");
-	}
-}
-
-void decudaInst::printPredicate()
-{
-	stringListPiece* currentPiece = m_predicate->getListStart();
-	if(currentPiece!=NULL)
-	{
-		output(currentPiece->stringText);
-
-		stringListPiece* currentPiece2 = m_predicateModifiers->getListStart();
-		for(int i=0; (i<m_predicateModifiers->getSize())&&(currentPiece2!=NULL); i++)
-		{
-			output(currentPiece2->stringText);
-
-			currentPiece2 = currentPiece2->nextString;
-		}
-
-		output(" ");
-	}
-}
-
-void decudaInst::printBaseModifiers()
-{
-	stringListPiece* currentPiece = m_baseModifiers->getListStart();
-	for(int i=0; (i<m_baseModifiers->getSize())&&(currentPiece!=NULL); i++)
-	{
-		output(currentPiece->stringText);
-
-		currentPiece = currentPiece->nextString;
-	}
-}
-
-void decudaInst::printTypeModifiers()
-{
-	stringListPiece* currentPiece = m_typeModifiers->getListStart();
-	for(int i=0; (i<m_typeModifiers->getSize())&&(currentPiece!=NULL); i++)
-	{
-		output(currentPiece->stringText);
-
-		currentPiece = currentPiece->nextString;
-	}
-}
-
-void decudaInst::printOperands()
-{
-	stringListPiece* currentPiece = m_operands->getListStart();
-	const char* operandDelimiter = "";
-	for(int i=0; (i<m_operands->getSize())&&(currentPiece!=NULL); i++)
-	{
-		output(operandDelimiter); output(" "); output(currentPiece->stringText);
-		operandDelimiter = ",";
-
-		currentPiece = currentPiece->nextString;
-	}
-}
-
-//This is where the conversion to new PTX takes place
-void decudaInst::printNewPtx()
-{
-	//
-	// Common modifications that apply to all instructions
-	//
-	stringListPiece* currentPiece;
-        int vectorFlag[4] = {0,0,0,0}; //0=16/32type, 1=bb64/ff64 type, 2=bb128 type
-
-	// Replace '%clock' with '%halfclock'
-	currentPiece = m_operands->getListStart();
-	for(int i=0; (i<m_operands->getSize())&&(currentPiece!=NULL); i++)
-	{
-		const char* modString = currentPiece->stringText;
-		if( strcmp(modString, "%%clock")==0 ) {
-			const char* newText = "%%halfclock";
-			currentPiece->stringText = newText;
-		}
-		currentPiece = currentPiece->nextString;
-	}
-
-	// Remove .join from base modifier list
-	currentPiece = m_baseModifiers->getListStart();
-	for(int i=0; (i<m_baseModifiers->getSize())&&(currentPiece!=NULL); i++)
-	{
-		const char* modString = currentPiece->stringText;
-		if( strcmp(modString, ".join")==0 ) {
-			m_baseModifiers->remove(i);
-		}
-		currentPiece = currentPiece->nextString;
-	}
-
-	// Change .end to .exit from base modified list
-	currentPiece = m_baseModifiers->getListStart();
-	for(int i=0; (i<m_baseModifiers->getSize())&&(currentPiece!=NULL); i++)
-	{
-		const char* modString = currentPiece->stringText;
-		if( strcmp(modString, ".end")==0 ) {
-			const char* newText = ".exit";
-			currentPiece->stringText = newText;
-		}
-		currentPiece = currentPiece->nextString;
-	}
-
-	// Change .b64 to .bb64 from type modifier list
-	// Change .b128 to .bb128 from type modifier list
-	// Change .f64 to .ff64 from type modifier list
-	currentPiece = m_typeModifiers->getListStart();
-	for(int i=0; (i<m_typeModifiers->getSize())&&(currentPiece!=NULL); i++)
-	{
-		const char* modString = currentPiece->stringText;
-		if( strcmp(modString, ".b64")==0 ) {
-			const char* newText = ".bb64";
-			currentPiece->stringText = newText;
-			vectorFlag[i] = 1;
-		}
-		else if( strcmp(modString, ".b128")==0 ) {
-			const char* newText = ".bb128";
-			currentPiece->stringText = newText;
-			vectorFlag[i] = 2;
-		}
-		else if( strcmp(modString, ".f64")==0 ) {
-			const char* newText = ".ff64";
-			currentPiece->stringText = newText;
-			vectorFlag[i] = 1;
-		}
-		currentPiece = currentPiece->nextString;
-	}
-
-	/*decuda bug workaround.
-	cvt.abs should drop the first type modifier.
-	cvt.rn.f32.f64 needs to have type modifiers reversed to cvt.ff64.f32*/
-	int absFound = 0;
-	if(strcmp(m_base, "cvt")==0 && m_typeModifiers->getSize() == 2)
-	{
-		stringListPiece* currentPieceCvt = m_baseModifiers->getListStart();
-		for(int i=0; (i<m_baseModifiers->getSize())&&(currentPieceCvt!=NULL); i++)
-		{
-                        const char* modStringCvt = currentPieceCvt->stringText;
-			if( strcmp(modStringCvt, ".abs")==0 ) {
-				vectorFlag[0] = vectorFlag[1];
-				absFound = 1;
-				break;
-			}
-			currentPieceCvt = currentPieceCvt->nextString;
-		}
-		if(absFound == 0 && vectorFlag[1] == 1)
-		{
-			vectorFlag[0] = 1;
-			vectorFlag[1] = 0;
-			const char* tempCharPtr = m_typeModifiers->getListStart()->stringText;
-			m_typeModifiers->getListStart()->stringText = m_typeModifiers->getListStart()->nextString->stringText;
-			m_typeModifiers->getListStart()->nextString->stringText = tempCharPtr;
-		}
-	}
-
-	/*decuda bug workaround, cvt.rz.f64 is really cvt.rz.f32.f64*/
-	if(strcmp(m_base, "cvt")==0 && m_typeModifiers->getSize() == 1 && vectorFlag[0]==1)
-	{
-		vectorFlag[0] = 0;
-		vectorFlag[1] = 1;
-		addTypeModifier(m_typeModifiers->getListStart()->stringText);
-		const char* newText = ".f32";
-		m_typeModifiers->getListStart()->stringText = newText;
-	}
-
-	/*expand vector operands eg. $r0 -> {$r0, $r1}*/
-	if((vectorFlag[0] != 0) || (vectorFlag[1] != 0) || (vectorFlag[2] != 0) || (vectorFlag[3] != 0) )
-	{
-		currentPiece = m_operands->getListStart();
-		for(int i=0; (i<m_operands->getSize())&&(currentPiece!=NULL); i++)
-		{
-			char *newText = new char[40];
-			char *regNumString;
-			int regNumInt;
-
-			const char* modString = currentPiece->stringText;
-
-			if(strcmp(m_base, "set")==0 || strcmp(m_base, "cvt")==0 || 
-				strcmp(m_base, "set?68?")==0 || strcmp(m_base, "set?65?")==0 ||
-				strcmp(m_base, "set?67?")==0 || strcmp(m_base, "set?13?")==0)
-			{
-				if( modString[0] == '$' && modString[1] == 'r' ) {
-					strcpy(newText, modString);
-					strtok (newText, "r");
-					regNumString = strtok (NULL, "r");
-					regNumInt = atoi(regNumString);
-					if(vectorFlag[i] ==0)
-						strcpy(newText, modString);
-					if(vectorFlag[i] ==1)
-						snprintf(newText,40,"{$r%u,$r%u}", regNumInt+0, regNumInt+1);
-					if(vectorFlag[i] ==2)
-						snprintf(newText,40,"{$r%u,$r%u,$r%u,$r%u}", regNumInt+0, regNumInt+1, regNumInt+2, regNumInt+3);
-
-					currentPiece->stringText = newText;
-				} else if( modString[0] == '-' && modString[1] == '$' && modString[2] == 'r' ) {
-					strcpy(newText, modString);
-					strtok (newText, "r");
-					regNumString = strtok (NULL, "r");
-					regNumInt = atoi(regNumString);
-					if(vectorFlag[i] ==0)
-						strcpy(newText, modString);
-					else if(vectorFlag[i] ==1)
-						snprintf(newText,40,"-{$r%u,$r%u}", regNumInt+0, regNumInt+1);
-					else if(vectorFlag[i] ==2)
-						snprintf(newText,40,"-{$r%u,$r%u,$r%u,$r%u}", regNumInt+0, regNumInt+1, regNumInt+2, regNumInt+3);
-
-					currentPiece->stringText = newText;
-				}
-			}
-			else
-			{
-				if( modString[0] == '$' && modString[1] == 'r' ) {
-					strcpy(newText, modString);
-					strtok (newText, "r");
-					regNumString = strtok (NULL, "r");
-					regNumInt = atoi(regNumString);
-					if(vectorFlag[0] ==0)
-						strcpy(newText, modString);
-					else if(vectorFlag[0] ==1)
-						snprintf(newText,40,"{$r%u,$r%u}", regNumInt+0, regNumInt+1);
-					else if(vectorFlag[0] ==2)
-						snprintf(newText,40,"{$r%u,$r%u,$r%u,$r%u}", regNumInt+0, regNumInt+1, regNumInt+2, regNumInt+3);
-
-					currentPiece->stringText = newText;
-				} else if( modString[0] == '-' && modString[1] == '$' && modString[2] == 'r' ) {
-					strcpy(newText, modString);
-					strtok (newText, "r");
-					regNumString = strtok (NULL, "r");
-					regNumInt = atoi(regNumString);
-					if(vectorFlag[0] ==0)
-						strcpy(newText, modString);
-					else if(vectorFlag[0] ==1)
-						snprintf(newText,40,"-{$r%u,$r%u}", regNumInt+0, regNumInt+1);
-					else if(vectorFlag[0] ==2)
-						snprintf(newText,40,"-{$r%u,$r%u,$r%u,$r%u}", regNumInt+0, regNumInt+1, regNumInt+2, regNumInt+3);
-
-					currentPiece->stringText = newText;
-				}
-			}
-			currentPiece = currentPiece->nextString;
-		}
-	}
-
-	//
-	// Instruction specific modifications
-	//
-	if(strcmp(m_base, "")==0)
-	{
-	}
-	else if(strcmp(m_base, ".entry")==0)
-	{
-		/*do nothing here*/
-	}
-	else if(strcmp(m_base, ".lmem")==0)
-	{
-	}
-	else if(strcmp(m_base, ".smem")==0)
-	{
-	}
-	else if(strcmp(m_base, ".reg")==0)
-	{
-	}
-	else if(strcmp(m_base, ".bar")==0)
-	{
-	}
-
-
-	else if(strcmp(m_base, "cvt")==0)
-	{
-
-		int cvt_inst_type = 0;	//0==cvt, 1==abs
-                int cvt_neg_mod_flag = 0; //0=off, 1=on
-
-		// Check the actual base instruction
-		stringListPiece* currentPiece = m_baseModifiers->getListStart();
-		for(int i=0; (i<m_baseModifiers->getSize())&&(currentPiece!=NULL); i++)
-		{
-                        const char* modString = currentPiece->stringText;
-			if( strcmp(modString, ".abs")==0 ) {
-				cvt_inst_type = 1;
-			}
-			if( strcmp(modString, ".neg")==0 ) {
-				cvt_neg_mod_flag = 1;
-			}
-			currentPiece = currentPiece->nextString;
-		}
-
-		if( cvt_inst_type == 1) {
-			// abs instruction
-			printLabel();
-			printPredicate();
-
-			output("abs");
-
-			if(m_typeModifiers->getSize() == 1)
-			{
-				const char* type = m_typeModifiers->getListStart()->stringText;
-				output(type);
-			}
-			else
-			{
-				const char* type = m_typeModifiers->getListStart()->nextString->stringText;
-				output(type);
-			}
-
-			printOperands();
-
-			output(";");
-
-		} else {
-			// cvt instruction
-			printLabel();
-			printPredicate();
-
-			output("cvt");
-
-			int typeModifiers = m_typeModifiers->getSize();
-
-			if(typeModifiers == 2) {
-
-				stringListPiece* currentPiece = m_baseModifiers->getListStart();
-
-				if(cvt_neg_mod_flag != 0)
-				    currentPiece = m_baseModifiers->getListStart()->nextString;
-
-				const char* dstType = m_typeModifiers->getListStart()->stringText;
-				const char* srcType = m_typeModifiers->getListStart()->nextString->stringText;
-
-				for(int i=0; (i<m_baseModifiers->getSize())&&(currentPiece!=NULL); i++)
-				{
-				        const char* modString = currentPiece->stringText;
-					if(
-						(strcmp(modString, ".rn")==0) || 
-						(strcmp(modString, ".rm")==0) ||
-						(strcmp(modString, ".rp")==0) ||
-						(strcmp(modString, ".rz")==0)
-					)
-					{
-						if((strcmp(dstType, ".f32")==0) && (strcmp(srcType, ".ff64")==0))
-							output(modString);
-						if((strcmp(dstType, ".f16")==0) && (strcmp(srcType, ".ff64")==0))
-							output(modString);
-						if((strcmp(dstType, ".f16")==0) && (strcmp(srcType, ".f32")==0))
-							output(modString);
-						if((strcmp(dstType, ".f16")==0) && (strcmp(srcType, ".u8")==0))
-							output(modString);
-						if((strcmp(dstType, ".f16")==0) && (strcmp(srcType, ".u16")==0))
-							output(modString);
-						if((strcmp(dstType, ".f16")==0) && (strcmp(srcType, ".u32")==0))
-							output(modString);
-						if((strcmp(dstType, ".f16")==0) && (strcmp(srcType, ".u64")==0))
-							output(modString);
-						if((strcmp(dstType, ".f16")==0) && (strcmp(srcType, ".s8")==0))
-							output(modString);
-						if((strcmp(dstType, ".f16")==0) && (strcmp(srcType, ".s16")==0))
-							output(modString);
-						if((strcmp(dstType, ".f16")==0) && (strcmp(srcType, ".s32")==0))
-							output(modString);
-						if((strcmp(dstType, ".f16")==0) && (strcmp(srcType, ".s64")==0))
-							output(modString);
-						if((strcmp(dstType, ".f32")==0) && (strcmp(srcType, ".u8")==0))
-							output(modString);
-						if((strcmp(dstType, ".f32")==0) && (strcmp(srcType, ".u16")==0))
-							output(modString);
-						if((strcmp(dstType, ".f32")==0) && (strcmp(srcType, ".u32")==0))
-							output(modString);
-						if((strcmp(dstType, ".f32")==0) && (strcmp(srcType, ".u64")==0))
-							output(modString);
-						if((strcmp(dstType, ".f32")==0) && (strcmp(srcType, ".s8")==0))
-							output(modString);
-						if((strcmp(dstType, ".f32")==0) && (strcmp(srcType, ".s16")==0))
-							output(modString);
-						if((strcmp(dstType, ".f32")==0) && (strcmp(srcType, ".s32")==0))
-							output(modString);
-						if((strcmp(dstType, ".f32")==0) && (strcmp(srcType, ".s64")==0))
-							output(modString);
-						if((strcmp(dstType, ".ff64")==0) && (strcmp(srcType, ".u8")==0))
-							output(modString);
-						if((strcmp(dstType, ".ff64")==0) && (strcmp(srcType, ".u16")==0))
-							output(modString);
-						if((strcmp(dstType, ".ff64")==0) && (strcmp(srcType, ".u32")==0))
-							output(modString);
-						if((strcmp(dstType, ".ff64")==0) && (strcmp(srcType, ".u64")==0))
-							output(modString);
-						if((strcmp(dstType, ".ff64")==0) && (strcmp(srcType, ".s8")==0))
-							output(modString);
-						if((strcmp(dstType, ".ff64")==0) && (strcmp(srcType, ".s16")==0))
-							output(modString);
-						if((strcmp(dstType, ".ff64")==0) && (strcmp(srcType, ".s32")==0))
-							output(modString);
-						if((strcmp(dstType, ".ff64")==0) && (strcmp(srcType, ".s64")==0))
-							output(modString);
-						if((strcmp(dstType, ".u8")==0) && (strcmp(srcType, ".f16")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".u16")==0) && (strcmp(srcType, ".f16")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".u32")==0) && (strcmp(srcType, ".f16")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".u64")==0) && (strcmp(srcType, ".f16")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".s8")==0) && (strcmp(srcType, ".f16")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".s16")==0) && (strcmp(srcType, ".f16")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".s32")==0) && (strcmp(srcType, ".f16")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".s64")==0) && (strcmp(srcType, ".f16")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".u8")==0) && (strcmp(srcType, ".f32")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".u16")==0) && (strcmp(srcType, ".f32")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".u32")==0) && (strcmp(srcType, ".f32")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".u64")==0) && (strcmp(srcType, ".f32")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".s8")==0) && (strcmp(srcType, ".f32")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".s16")==0) && (strcmp(srcType, ".f32")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".s32")==0) && (strcmp(srcType, ".f32")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".s64")==0) && (strcmp(srcType, ".f32")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".u8")==0) && (strcmp(srcType, ".ff64")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".u16")==0) && (strcmp(srcType, ".ff64")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".u32")==0) && (strcmp(srcType, ".ff64")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".u64")==0) && (strcmp(srcType, ".ff64")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".s8")==0) && (strcmp(srcType, ".ff64")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".s16")==0) && (strcmp(srcType, ".ff64")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".s32")==0) && (strcmp(srcType, ".ff64")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".s64")==0) && (strcmp(srcType, ".ff64")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".f16")==0) && (strcmp(srcType, ".f16")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".f32")==0) && (strcmp(srcType, ".f32")==0))
-							{ output(modString); output("i"); }
-						if((strcmp(dstType, ".ff64")==0) && (strcmp(srcType, ".ff64")==0))
-							{ output(modString); output("i"); }
-					}
-					else if (						
-						(strcmp(modString, ".rni")==0 || strcmp(modString, ".rmi")==0 ||
-						 strcmp(modString, ".rpi")==0 || strcmp(modString, ".rzi")==0)
-					)
-					{
-						output(modString);
-					}
-					else
-					{
-						output("\nUnknown mod:"); output(currentPiece->stringText);
-						assert(0);
-					}
-
-					currentPiece = currentPiece->nextString;
-				}
-			} else {
-				printBaseModifiers();
-			}
-
-			// If one type modifier, duplicate to two;
-			int numModifiers = m_typeModifiers->getSize();
-			stringListPiece* currentPiece = m_typeModifiers->getListStart();
-			if( numModifiers == 1 ) {
-				output(currentPiece->stringText);
-				output(currentPiece->stringText);
-			} else {
-				printTypeModifiers();
-			}
-
-			printOperands();
-
-			output(";");
-
-		}
-		
-	}
-	else if(strcmp(m_base, "shl")==0)
-	{
-		printDefaultPtx();
-	}
-	else if(strcmp(m_base, "add")==0)
-	{
-		printLabel();
-		printPredicate();
-
-		output(m_base);
-
-		printBaseModifiers();
-
-		currentPiece = m_typeModifiers->getListStart();
-		for(int i=0; (i<m_typeModifiers->getSize())&&(currentPiece!=NULL); i++)
-		{
-			if(strcmp(currentPiece->stringText, ".b8")==0)
-			{
-				output(".u8");
-			}
-			else if(strcmp(currentPiece->stringText, ".b16")==0)
-			{
-				output(".u16");
-			}
-			else if(strcmp(currentPiece->stringText, ".b32")==0)
-			{
-				output(".u32");
-			}
-			else if(strcmp(currentPiece->stringText, ".b64")==0)
-			{
-				output(".u64");
-			}
-			else
-			{
-				output(currentPiece->stringText);
-			}
-
-			currentPiece = currentPiece->nextString;
-		}
-
-		printOperands();
-
-		output(";");
-	}
-	else if(strcmp(m_base, "movsh")==0)
-	{
-		printLabel();
-		printPredicate();
-
-		output("shl");
-
-		printBaseModifiers();
-
-		printTypeModifiers();
-
-		printOperands();
-
-		output(";");
-	}
-	else if(strcmp(m_base, "mov")==0)
-	{
-		printLabel();
-		printPredicate();
-
-		output("mov");
-
-		printBaseModifiers();
-
-		// If mov has two type modifiers, pick the smaller one for proper
-		// bit truncation
-		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;
-			output(type.c_str());
-		} else if( m_typeModifiers->getSize() == 1 ) {
-			printTypeModifiers();
-		} else {
-			output("Error: unsupported number of type modifiers. ");
-		}
-
-		printOperands();
-
-		output(";");
-	}
-	else if(strcmp(m_base, "bar.sync")==0)
-	{
-		printLabel();
-		printPredicate();
-
-		output(m_base);
-
-		printBaseModifiers();
-
-		stringListPiece* currentPiece = m_typeModifiers->getListStart();
-		currentPiece = m_typeModifiers->getListStart();
-		for(int i=0; (i<m_typeModifiers->getSize())&&(currentPiece!=NULL); i++)
-		{
-			//output(currentPiece->stringText);
-
-			currentPiece = currentPiece->nextString;
-		}
-
-		printOperands();
-
-		output(";");
-      m_opPerCycle = -1;
-	}
-	else if(strcmp(m_base, "mul")==0)
-	{
-		printDefaultPtx();
-
-      // opPerCycle - lower if a 32-bit integer mul
-      const char* dstType = m_typeModifiers->getListStart()->stringText;
-      if( strcmp(dstType, ".s32")==0 || strcmp(dstType, ".u32")==0 || strcmp(dstType, ".b32")==0 )
-         m_opPerCycle = 2;
-	}
-	else if(strcmp(m_base, "mad24")==0)
-	{
-		printLabel();
-		printPredicate();
-
-		output("mad24");
-
-		printBaseModifiers();
-
-		// Only output the destination operand type (first type modifier only)
-		stringListPiece* currentPiece = m_typeModifiers->getListStart();
-		output(currentPiece->stringText);
-
-		printOperands();
-		// If 3 operands, this is a mac24 instruction, so add destination operand as 4th operand
-		int numOperands = m_operands->getSize();
-		if(numOperands == 3) {
-			stringListPiece* currentPiece = m_operands->getListStart();
-			output(", "); output(currentPiece->stringText);
-		}
-
-		output(";");
-	}
-	else if(strcmp(m_base, "mad24c1")==0)
-	{
-		output("nop; //");
-
-		printLabel();
-		printPredicate();
-
-		output("mad24c1");
-
-		printBaseModifiers();
-
-		// Only output the destination operand type (first type modifier only)
-		stringListPiece* currentPiece = m_typeModifiers->getListStart();
-		output(currentPiece->stringText);
-
-		printOperands();
-		// If 3 operands, this is a mac24 instruction, so add destination operand as 4th operand
-		int numOperands = m_operands->getSize();
-		if(numOperands == 3) {
-			stringListPiece* currentPiece = m_operands->getListStart();
-			output(", "); output(currentPiece->stringText);
-		}
-
-		output(";");
-	}
-	else if(strcmp(m_base, "set")==0)
-	{
-		printLabel();
-		printPredicate();
-
-		output("set");
-
-		printBaseModifiers();
-
-		// If one type modifier, duplicate to two;
-		// if three type modifiers, remove third
-		int numModifiers = m_typeModifiers->getSize();
-		stringListPiece* currentPiece = m_typeModifiers->getListStart();
-		if( numModifiers == 1 ) {
-			output(currentPiece->stringText);
-			output(currentPiece->stringText);
-		} else if( numModifiers == 3) {
-			for(int i=0; (i<numModifiers-1)&&(currentPiece!=NULL); i++)
-			{
-				output(currentPiece->stringText);
-				currentPiece = currentPiece->nextString;
-			}
-		}
-
-		printOperands();
-
-		output(";");
-	}
-	else if(strcmp(m_base, "mad")==0)
-	{
-		printDefaultPtx();
-
-      // opPerCycle - lower if a 32-bit integer mad
-      const char* dstType = m_typeModifiers->getListStart()->stringText;
-      if( strcmp(dstType, ".s32")==0 || strcmp(dstType, ".u32")==0 || strcmp(dstType, ".b32")==0 )
-         m_opPerCycle = 2;
-	}
-	else if(strcmp(m_base, "mul24")==0)
-	{
-		printLabel();
-		printPredicate();
-
-		output("mul24");
-
-		printBaseModifiers();
-
-		// Only output the destination operand type (first type modifier only)
-		stringListPiece* currentPiece = m_typeModifiers->getListStart();
-		output(currentPiece->stringText);
-
-		printOperands();
-
-		output(";");
-	}
-	else if(strcmp(m_base, "set?68?")==0)
-	{
-		// This actually takes absolute value of first source operand and is set.gt
-		printLabel();
-		printPredicate();
-
-		output("set.gt.abs");
-
-		printBaseModifiers();
-
-		// If one type modifier, duplicate to two;
-		// if three type modifiers, remove third
-		int numModifiers = m_typeModifiers->getSize();
-		stringListPiece* currentPiece = m_typeModifiers->getListStart();
-		if( numModifiers == 1 ) {
-			output(currentPiece->stringText);
-			output(currentPiece->stringText);
-		} else if( numModifiers == 3) {
-			for(int i=0; (i<numModifiers-1)&&(currentPiece!=NULL); i++)
-			{
-				output(currentPiece->stringText);
-				currentPiece = currentPiece->nextString;
-			}
-		}
-
-		printOperands();
-
-		output(";");
-
-		output(" //set?68?");
-	}
-	else if(strcmp(m_base, "set?65?")==0)
-	{
-		// This actually takes absolute value of first source operand and is set.lt
-		printLabel();
-		printPredicate();
-
-		output("set.lt.abs");
-
-		printBaseModifiers();
-
-		// If one type modifier, duplicate to two;
-		// if three type modifiers, remove third
-		int numModifiers = m_typeModifiers->getSize();
-		stringListPiece* currentPiece = m_typeModifiers->getListStart();
-		if( numModifiers == 1 ) {
-			output(currentPiece->stringText);
-			output(currentPiece->stringText);
-		} else if( numModifiers == 3) {
-			for(int i=0; (i<numModifiers-1)&&(currentPiece!=NULL); i++)
-			{
-				output(currentPiece->stringText);
-				currentPiece = currentPiece->nextString;
-			}
-		}
-
-		printOperands();
-
-		output(";");
-
-		output(" //set?65?");
-	}
-	else if(strcmp(m_base, "set?67?")==0)
-	{
-		// Change to set.gt
-		printLabel();
-		printPredicate();
-
-		output("set.gt.abs");
-
-		printBaseModifiers();
-
-		// If one type modifier, duplicate to two;
-		// if three type modifiers, remove third
-		int numModifiers = m_typeModifiers->getSize();
-		stringListPiece* currentPiece = m_typeModifiers->getListStart();
-		if( numModifiers == 1 ) {
-			output(currentPiece->stringText);
-			output(currentPiece->stringText);
-		} else if( numModifiers == 3) {
-			for(int i=0; (i<numModifiers-1)&&(currentPiece!=NULL); i++)
-			{
-				output(currentPiece->stringText);
-				currentPiece = currentPiece->nextString;
-			}
-		}
-
-		printOperands();
-
-		output(";");
-	}
-	else if(strcmp(m_base, "set?13?")==0)
-	{
-		// Change to set.gt
-		printLabel();
-		printPredicate();
-
-		output("set.neu");
-
-		printBaseModifiers();
-
-		// If one type modifier, duplicate to two;
-		// if three type modifiers, remove third
-		int numModifiers = m_typeModifiers->getSize();
-		stringListPiece* currentPiece = m_typeModifiers->getListStart();
-		if( numModifiers == 1 ) {
-			output(currentPiece->stringText);
-			output(currentPiece->stringText);
-		} else if( numModifiers == 3) {
-			for(int i=0; (i<numModifiers-1)&&(currentPiece!=NULL); i++)
-			{
-				output(currentPiece->stringText);
-				currentPiece = currentPiece->nextString;
-			}
-		}
-
-		printOperands();
-
-		output(";");
-	}
-	else if(strcmp(m_base, "rcp")==0)
-	{
-		printDefaultPtx();
-
-      m_opPerCycle = 2;
-	}
-	else if(strcmp(m_base, "pre.sin")==0)
-	{
-		printLabel();
-		printPredicate();
-
-		output("nop;");
-		output(" //");
-		output("pre.sin");
-		printBaseModifiers();
-		printTypeModifiers();
-		printOperands();
-
-      // m_opPerCycle - Ignore
-      m_opPerCycle = -1;
-	}
-	else if(strcmp(m_base, "sin")==0)
-	{
-		printDefaultPtx();
-
-      m_opPerCycle = 1;
-	}
-	else if(strcmp(m_base, "pre.ex2")==0)
-	{
-		printLabel();
-		printPredicate();
-
-		output("ex2");
-		printBaseModifiers();
-		printTypeModifiers();
-		printOperands();
-		output(";");
-
-		output(" //");
-		printDefaultPtx();
-
-      // m_opPerCycle - Ignore
-      m_opPerCycle = -1;
-	}
-	else if(strcmp(m_base, "ex2")==0)
-	{
-		printLabel();
-		printPredicate();
-
-		output("nop;");
-		output(" //");
-		output("ex2");
-		printBaseModifiers();
-		printTypeModifiers();
-		printOperands();
-
-      m_opPerCycle = 1;
-	}
-	else if(strcmp(m_base, "cos")==0)
-	{
-		printDefaultPtx();
-
-      m_opPerCycle = 1;
-	}
-	else if(strcmp(m_base, "lg2")==0)
-	{
-		printDefaultPtx();
-
-      m_opPerCycle = 2;
-	}
-	else if(strcmp(m_base, "rsqrt")==0)
-	{
-		printDefaultPtx();
-
-      m_opPerCycle = 2;
-	}
-	else if(strcmp(m_base, "mac")==0)
-	{
-		// Replace with mad by adding a 4th operand
-		printLabel();
-		printPredicate();
-
-		output("mad");
-
-		printBaseModifiers();
-
-		printTypeModifiers();
-
-		// Print operands and then include destination (1st) operand as the 4th operand
-		printOperands();
-		stringListPiece* currentPiece = m_operands->getListStart();
-		output(","); output(" "); output(currentPiece->stringText);
-
-		output(";");
-
-      // opPerCycle - lower if a 32-bit integer mac
-      const char* dstType = m_typeModifiers->getListStart()->stringText;
-      if( strcmp(dstType, ".s32")==0 || strcmp(dstType, ".u32")==0 || strcmp(dstType, ".b32")==0 )
-         m_opPerCycle = 2;
-	}
-	else if(strcmp(m_base, "bra.label")==0)
-	{
-		printLabel();
-		printPredicate();
-
-		output("bra");
-
-		printBaseModifiers();
-
-		printTypeModifiers();
-
-		printOperands();
-
-		output(";");
-	}
-	else if(strcmp(m_base, "join.label")==0)
-	{
-		printLabel();
-		printPredicate();
-
-		output("nop;");
-		output(" //join.label");
-
-      // Ignore
-      m_opPerCycle = -1;
-	}
-	else if(strcmp(m_base, "nop.end")==0)
-	{
-		printLabel();
-		printPredicate();
-
-		output("nop.exit;");
-
-      // Ignore
-      m_opPerCycle = -1;
-	}
-	else if(strcmp(m_base, "nop.join")==0)
-	{
-		printLabel();
-		printPredicate();
-
-		output("nop;");
-
-      // Ignore
-      m_opPerCycle = -1;
-	}
-	else if(strcmp(m_base, "nop")==0)
-	{
-		//printLabel();
-                //printPredicate();
-
-                output("nop;");
-	}
-	else if(strcmp(m_base, "return")==0)
-	{
-		// ret instruction causes a deadlock bug in the simulator
-		// Temporary fix: branch to a dummy exit instruction that is added to the end of each
-		//                entry with the label 'l_exit'
-		printLabel();
-		printPredicate();
-
-		output("retp");
-		//output("bra l_exit");
-
-		printBaseModifiers();
-
-		printTypeModifiers();
-
-		printOperands();
-
-		output(";");
-	}
-	else if(strcmp(m_base, "and")==0)
-	{
-		printDefaultPtx();
-	}
-	else if(strcmp(m_base, "andn")==0)
-	{
-		printDefaultPtx();
-	}
-	else if(strcmp(m_base, "tex")==0)
-	{
-		printDefaultPtx();
-	}
-	else if(strcmp(m_base, "xor")==0)
-	{
-		printDefaultPtx();
-	}
-	else if(strcmp(m_base, "or")==0)
-	{
-		printDefaultPtx();
-	}
-	else if(strcmp(m_base, "shr")==0)
-	{
-		printDefaultPtx();
-	}
-	else if(strcmp(m_base, "subr")==0)
-	{
-
-		// Replace with sub instruction with destination operands switched
-		printLabel();
-		printPredicate();
-
-		output("sub");
-
-		printBaseModifiers();
-
-		printTypeModifiers();
-
-		// Switch source operands before printing
-		// Must be 3 operands, switch the last two;
-		if( m_operands->getSize() != 3) {
-			output("Error: subr instruction with number of operands other than 3.\n");
-			assert(0);
-		}
-		const char* firstOperand = m_operands->getListStart()->stringText;
-		const char* secondOperand = m_operands->getListStart()->nextString->stringText;
-		const char* thirdOperand = m_operands->getListStart()->nextString->nextString->stringText;
-		output(" "); output(firstOperand); output(",");
-		output(" "); output(thirdOperand); output(",");
-		output(" "); output(secondOperand);
-		
-
-		output(";");
-	}
-	else if(strcmp(m_base, "sub")==0)
-	{
-		printDefaultPtx();
-	}
-	else if(strcmp(m_base, "max")==0)
-	{
-		printDefaultPtx();
-	}
-	else if(strcmp(m_base, "min")==0)
-	{
-		printDefaultPtx();
-	}
-	else if(strcmp(m_base, "call.label")==0)
-	{
-		printLabel();
-		printPredicate();
-
-		output("callp");
-
-		printBaseModifiers();
-
-		printTypeModifiers();
-
-		printOperands();
-
-		output(";");
-
-      // Ignore
-      m_opPerCycle = -1;
-	}
-	else if(strcmp(m_base, "not")==0)
-	{
-		printDefaultPtx();
-	}
-	else if(strcmp(m_base, "delta")==0)
-	{
-		// This is a neg instruction
-		printLabel();
-		printPredicate();
-
-		output("neg");
-
-		printBaseModifiers();
-
-		printTypeModifiers();
-
-		printOperands();
-
-		output(";");
-	}
-	else if(strcmp(m_base, "break")==0)
-	{
-		printDefaultPtx();
-	}
-	else if(strcmp(m_base, "breakaddr.label")==0)
-	{
-		printLabel();
-		printPredicate();
-
-		output("breakaddr");
-
-		printBaseModifiers();
-
-		printTypeModifiers();
-
-		printOperands();
-
-		output(";");
-	}
-	else if(strcmp(m_base, "inc")==0)
-	{
-		printDefaultPtx();
-	}
-	else if(strcmp(m_base, "exch")==0)
-	{
-		printDefaultPtx();
-	}
-	else if(strcmp(m_base, "cas")==0)
-	{
-		printDefaultPtx();
-	}
-   	else if(strcmp(m_base, "norn")==0)
-	{
-		printDefaultPtx();
-	}
-	else if(strcmp(m_base, "addc")==0)
-	{
-		output("nop; //");
-		printDefaultPtx();
-	}
-	else if(strcmp(m_base, "orn")==0)
-	{
-		printDefaultPtx();
-	}
-	else if(strcmp(m_base, "nandn")==0)
-	{
-		printDefaultPtx();
-	}
-   	else if(strcmp(m_base, "nxor")==0)
-	{
-		output("nop; //");
-		printDefaultPtx();
-	}
-   	else if(strcmp(m_base, "sad")==0)
-	{
-		output("nop; //");
-		printDefaultPtx();
-	}
-   	else if(strcmp(m_base, "op.13")==0)
-	{
-		output("nop; //");
-		printDefaultPtx();
-	}
-   	else if(strcmp(m_base, "op.e5")==0)
-	{
-		output("nop; //");
-		printDefaultPtx();
-	}
-   	else if(strcmp(m_base, "op.e6")==0)
-	{
-		output("nop; //");
-		printDefaultPtx();
-	}
-   	else if(strcmp(m_base, "op.d0")==0)
-	{
-		output("nop; //");
-		printDefaultPtx();
-	}
-
-
-	else if(strcmp(m_base, "{")==0)
-	{
-		//output(m_base);
-
-      // Ignore
-      m_opPerCycle = -1;
-	}
-	else if(strcmp(m_base, "}")==0)
-	{
-		//output(m_base);
-
-      // Ignore
-      m_opPerCycle = -1;
-	}
-	else
-	{
-		output("Unknown Instruction: "); output(m_base);
-		assert(0);
-	}
-
-}
diff --git a/decuda_to_ptxplus/decudaInst.h b/decuda_to_ptxplus/decudaInst.h
index 7011130..491fe3c 100644
--- a/decuda_to_ptxplus/decudaInst.h
+++ b/decuda_to_ptxplus/decudaInst.h
@@ -25,7 +25,8 @@
 // OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
 // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
-
+#ifndef _DECUDAINST_H_
+#define _DECUDAINST_H_
 #include "stringList.h"
 #include <assert.h>
 
@@ -92,3 +93,4 @@ public:
 	//TODO: translate to New PTX and print
 
 };
+#endif //_DECUDAINST_H_
diff --git a/decuda_to_ptxplus/decudaInstList.cc b/decuda_to_ptxplus/decudaInstList.cc
deleted file mode 100644
index 4035d76..0000000
--- a/decuda_to_ptxplus/decudaInstList.cc
+++ /dev/null
@@ -1,836 +0,0 @@
-// Copyright (c) 2009-2011, Jimmy Kwa,
-// The University of British Columbia
-// All rights reserved.
-//
-// Redistribution and use in source and binary forms, with or without
-// modification, are permitted provided that the following conditions are met:
-//
-// Redistributions of source code must retain the above copyright notice, this
-// list of conditions and the following disclaimer.
-// Redistributions in binary form must reproduce the above copyright notice, this
-// list of conditions and the following disclaimer in the documentation and/or
-// other materials provided with the distribution.
-// Neither the name of The University of British Columbia nor the names of its
-// contributors may be used to endorse or promote products derived from this
-// software without specific prior written permission.
-//
-// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
-// ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
-// WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
-// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
-// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
-// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
-// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
-// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
-// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
-// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-
-#include "decudaInstList.h"
-#define P_DEBUG 1
-#define DPRINTF(...) \
-   if(P_DEBUG) { \
-      printf("(%s:%u) ", __FILE__, __LINE__); \
-      printf(__VA_ARGS__); \
-      printf("\n"); \
-      fflush(stdout); \
-   }
-extern void output(const char * text);
-
-//Constructor
-decudaInstList::decudaInstList()
-{
-	//initilize everything to empty
-}
-
-//retrieve point to list end
-decudaInst decudaInstList::getListEnd()
-{
-	return m_entryList.back().m_instList.back();
-}
-
-//add decudaInst to the last entry in entry list
-int decudaInstList::add(decudaInst* newDecudaInst)
-{
-	if(m_entryList.size() == 0) {
-		//output("ERROR: Adding an instruction before entry.\n");
-		addEntry("");
-		//assert(0);
-	}
-
-	m_entryList.back().m_instList.push_back(*newDecudaInst);
-
-	return m_entryList.size();
-}
-
-// add a new entry
-int decudaInstList::addEntry(std::string entryName)
-{
-	decudaEntry newEntry;
-	newEntry.m_largestRegIndex = -1;
-	newEntry.m_largestOfsRegIndex = -1;
-	newEntry.m_largestPredIndex = -1;
-	newEntry.m_reg124 = false;
-	newEntry.m_oreg127 = false;
-	newEntry.m_lMemSize = -1;
-
-	newEntry.m_entryName = entryName;
-
-
-   // Fill opPerCycle histogram with values
-   newEntry.m_opPerCycleHistogram.insert( std::pair<std::string,int>("OP_1", 0) );
-   newEntry.m_opPerCycleHistogram.insert( std::pair<std::string,int>("OP_2", 0) );
-   newEntry.m_opPerCycleHistogram.insert( std::pair<std::string,int>("OP_8", 0) );
-
-
-	m_entryList.push_back(newEntry);
-	return m_entryList.size();
-}
-
-void decudaInstList::setLastEntryName(std::string entryName)
-{
-	m_entryList.back().m_entryName = entryName;
-}
-
-void decudaInstList::setLastEntryLMemSize(int lMemSize)
-{
-	m_entryList.back().m_lMemSize = lMemSize;
-}
-
-bool decudaInstList::findEntry(std::string entryName, decudaEntry& entry) {
-	std::list<decudaEntry>::iterator e;
-
-	std::string entryNameS = entryName;
-
-	for(e=m_entryList.begin(); e!=m_entryList.end(); ++e) {
-		if( e->m_entryName == entryNameS) {
-			entry = *e;
-			return true;
-		}
-	}
-
-	return false;
-}
-
-void decudaInstList::printEntryNames() {
-	printf("------------\n");
-	printf("%d Entry names:\n", m_entryList.size());
-	std::list<decudaEntry>::iterator e;
-	for(e=m_entryList.begin(); e!=m_entryList.end(); ++e) {
-		printf("existing entry=%s\n", e->m_entryName.c_str());
-	}
-	printf("------------\n");
-}
-
-// print out .version and .target headers
-void decudaInstList::printHeaderInstList()
-{
-	// These should be in the first entry
-	decudaEntry e_first = m_entryList.front();
-
-	std::list<decudaInst>::iterator currentInst;
-	for(currentInst=e_first.m_instList.begin(); currentInst!=e_first.m_instList.end(); ++currentInst)
-	{
-		if(!(currentInst->printHeaderInst()))
-		{
-			break;
-		}
-	}
-	for (	std::list<std::string>::iterator iter = m_realTexList.begin();
-			iter != m_realTexList.end();
-			iter ++) {
-		output(".tex .u64 ");
-		output((*iter).c_str());
-		output(";\n");
-	}
-}
-
-void decudaInstList::printNewPtxList(decudaInstList* headerInfo)
-{
-
-	// Print memory segment definitions
-	output("\n");
-	printMemory();
-
-
-	// Each entry
-	std::list<decudaEntry>::iterator e;
-	for(e=m_entryList.begin(); e!=m_entryList.end(); ++e) {
-
-		// The first instruction will be the entry instruction
-		std::list<decudaInst>::iterator currentInst;
-		currentInst=e->m_instList.begin();
-
-		// 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 headerInst;
-			for(headerInst=headerEntry.m_instList.begin(); 
-			     headerInst!=headerEntry.m_instList.end(); 
-			     ++headerInst) {
-				if(headerInst!=headerEntry.m_instList.begin())
-					output("\t");
-				headerInst->printHeaderPtx();
-				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");
-			} else {
-				output("Mismatch in entry names between decuda output and original ptx file.\n");
-				assert(0);
-			}
-		}
-
-		// Output the registers, predicates and other things
-		output("{\n");
-		printRegNames(*e);
-		printPredNames(*e);
-		printOutOfBoundRegisters(*e);
-		output("\n");
-
-
-		// Print the rest of the instructions in this entry
-		for(++currentInst; currentInst!=e->m_instList.end(); ++currentInst){
-			// Output the instruction
-			output("\t");
-			currentInst->printNewPtx();
-			output("\n");
-
-         // Update the opPerCycle histogram
-         int opPerCycle = currentInst->getOpPerCycle();
-         switch( opPerCycle ) {
-            case 8:
-               e->m_opPerCycleHistogram["OP_8"] += 1;
-               break;
-            case 2:
-               e->m_opPerCycleHistogram["OP_2"] += 1;
-               break;
-            case 1:
-               e->m_opPerCycleHistogram["OP_1"] += 1;
-               break;
-         }
-		}
-
-		// To prevent the 'ret' instruction deadlock bug in gpgpusim, insert a dummy exit instruction
-		output("\n\t");
-		output("l_exit: exit;");
-		output("\n");
-
-		output("}\n\n\n");
-
-      // Print out histogram
-      printf("Entry: %s\n", e->m_entryName.c_str());
-      printf("OP_8 %d\n", e->m_opPerCycleHistogram["OP_8"]);
-      printf("OP_2 %d\n", e->m_opPerCycleHistogram["OP_2"]);
-      printf("OP_1 %d\n", e->m_opPerCycleHistogram["OP_1"]);
-      printf("\n");
-	}
-}
-
-
-// print out register names
-void decudaInstList::printRegNames(decudaEntry entry)
-{
-	if( entry.m_largestRegIndex >= 0) {
-		char out[30];
-		sprintf(out, "\t.reg .u32 $r<%d>;", entry.m_largestRegIndex+1);
-		output(out);
-		output("\n");
-	}
-
-	if( entry.m_largestOfsRegIndex >= 0) {
-		char out[30];
-		sprintf(out, "\t.reg .u32 $ofs<%d>;", entry.m_largestOfsRegIndex+1);
-		output(out);
-		output("\n");
-	}
-}
-
-// print reg124 and set its value to 0
-void decudaInstList::printOutOfBoundRegisters(decudaEntry entry)
-{
-	if( entry.m_reg124 == true ) {
-		output("\n");
-		output("\t.reg .u32 $r124;\n");
-		output("\tmov.u32 $r124, 0x00000000;\n");
-	}
-	if( entry.m_oreg127 == true) {
-		output("\n");
-		output("\t.reg .u32 $o127;\n");
-	}
-}
-
-// increment register list and parse register
-std::string decudaInstList::parseRegister(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 '.lo' or '.hi')
-	if(lo)
-		reg = reg.substr(0, reg.size()-3);
-
-
-	// Increase register number if needed
-	// Two types of registers, $r# or $ofs#
-	if(reg.substr(0, 2) == "$r") {
-		reg = reg.substr(2, reg.size()-2);
-		int regNum = atoi(reg.c_str());
-
-		// Remove register overlap at 64
-		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, 4) == "$ofs") {
-		reg = reg.substr(4, reg.size()-4);
-		int regNum = atoi(reg.c_str());
-
-		if( m_entryList.back().m_largestOfsRegIndex < regNum && regNum < 124 )
-			m_entryList.back().m_largestOfsRegIndex = regNum;
-	} else if(reg == "$o127") {
-		m_entryList.back().m_oreg127 = true;
-	} else {
-		output("ERROR: unknown register type.\n");
-		assert(0);
-	}
-	return origReg;
-}
-
-
-// add to register list
-void decudaInstList::addRegister(std::string reg, bool lo)
-{
-	//Check to see if the register is an implied vector.
-	//If .b64 is a type modifier, $r0 becomes {$r0, $r1}
-	//If .b128 is a type modifier, $r0 becomes {$r0, $r1, $r2, $r3}
-	//This information is passed to parseRegister so the registers get declared.
-	int vectorFlag = 0;
-	stringList* typeModifiers = getListEnd().getTypeModifiers();
-	stringListPiece* currentPiece;
-	currentPiece = typeModifiers->getListStart();
-        for(int i=0; (i<typeModifiers->getSize())&&(currentPiece!=NULL); i++)
-        {
-                const char* modString = currentPiece->stringText;
-
-                if( (strcmp(modString, ".b64")==0) || (strcmp(modString, ".f64")==0) )
-                        vectorFlag = 64;
-                if( strcmp(modString, ".b128")==0 )
-                        vectorFlag = 128;
-
-                currentPiece = currentPiece->nextString;
-        }
-
-
-	std::string parsedReg = parseRegister(reg, lo, vectorFlag);
-
-	// Add the register to instruction operand list
-	char* regName = new char [strlen(parsedReg.c_str())+1];
-	strcpy(regName, parsedReg.c_str());
-	getListEnd().addOperand(regName);
-}
-
-
-
-// print out predicate names
-void decudaInstList::printPredNames(decudaEntry entry)
-{
-	if( entry.m_largestPredIndex >= 0) {
-		char out[30];
-		sprintf(out, "\t.reg .pred $p<%d>;", entry.m_largestPredIndex+1);
-		output(out);
-		output("\n");
-	}
-	
-}
-
-// increment predicate list
-std::string decudaInstList::parsePredicate(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;
-}
-
-// add to predicate list
-void decudaInstList::addPredicate(std::string pred)
-{
-	std::string parsedPred = parsePredicate(pred);
-
-	// Add the predicate to instruction operand list
-	char* predName = new char [strlen(parsedPred.c_str())+1];
-	strcpy(predName, parsedPred.c_str());
-	getListEnd().addOperand(predName);
-}
-
-
-// pred|reg double operand
-void decudaInstList::addDoublePredReg(std::string pred, std::string reg, bool lo)
-{
-	std::string parsedPred = parsePredicate(pred);
-	std::string parsedReg = parseRegister(reg, lo, 0);
-
-	// Add the double operand to instruction operand list
-	// If the base instruction is "set", then both operand get same value, use '/' for separator
-	// For cvt,shr,mul use '|' separator
-	std::string doublePredReg;
-	if( 
-		strcmp(getListEnd().getBase(), "set") == 0 ||
-		strcmp(getListEnd().getBase(), "setp") == 0 ||
-		strcmp(getListEnd().getBase(), "set?68?") == 0 ||
-		strcmp(getListEnd().getBase(), "set?65?") == 0 ||
-		strcmp(getListEnd().getBase(), "set?67?") == 0 ||
-		strcmp(getListEnd().getBase(), "set?13?") == 0
-	)
-		doublePredReg = parsedPred + "/" + parsedReg;
-	else
-		doublePredReg = parsedPred + "|" + parsedReg;
-
-	char* doublePredRegName = new char [strlen(doublePredReg.c_str())+1];
-	strcpy(doublePredRegName, doublePredReg.c_str());
-	getListEnd().addOperand(doublePredRegName);
-}
-
-
-// add to tex list
-void decudaInstList::addTex(std::string tex)
-{
-	std::string origTex = tex;
-	DPRINTF("decudaInstList::addTex tex=%s", tex.c_str());
-	// If $tex# tex from decuda, then use index to get real tex name
-	if(tex.substr(0, 4) == "$tex") {
-		tex = tex.substr(4, tex.size()-4);
-		int texNum = atoi(tex.c_str());
-		if(texNum >= m_realTexList.size()) {
-			output("ERROR: tex does not exist in real tex list from ptx.\n.");
-			assert(0);
-		}
-		
-		std::list<std::string>::iterator itex = m_realTexList.begin();
-		for(int i=0; i<texNum; i++) itex++;
-		origTex = *itex;
-	}
-	// Otherwise, tex from original ptx
-	else {
-		m_realTexList.push_back(tex);
-	}
-
-	// Add the tex to instruction operand list
-	//char* texName = new char [strlen(origTex.c_str())+1];
-	//strcpy(texName, origTex.c_str());
-	//getListEnd().addOperand(texName);
-}
-
-
-// create new global constant memory "bank"
-void decudaInstList::addConstMemory(int index)
-{
-	constMemory newConstMem;
-	newConstMem.index = index;
-	newConstMem.entryIndex = 0;
-	m_constMemoryList.push_back(newConstMem);
-}
-
-// create new entry specific constant memory "bank"
-void decudaInstList::addEntryConstMemory(int index)
-{
-	constMemory newConstMem;
-	newConstMem.index = index;
-	newConstMem.entryIndex = m_entryList.size();
-	m_constMemoryList.push_back(newConstMem);
-}
-
-
-
-// add value to const memory
-void decudaInstList::addConstMemoryValue(std::string constMemoryValue)
-{
-	m_constMemoryList.back().m_constMemory.push_back(constMemoryValue);
-}
-
-void decudaInstList::addConstMemoryValue2(std::string constMemoryValue)
-{
-	m_constMemoryList2.back().m_constMemory.push_back(constMemoryValue);
-}
-
-// set type of constant memory
-void decudaInstList::setConstMemoryType(const char* type)
-{
-	m_constMemoryList.back().type = type;
-}
-
-void decudaInstList::setConstMemoryType2(const char* type)
-{
-	m_constMemoryList2.back().type = type;
-}
-
-// print const memory directive
-void decudaInstList::printMemory()
-{
-
-	// Constant memory
-
-	for(std::list<constMemory>::iterator i=m_constMemoryList.begin(); i!=m_constMemoryList.end(); ++i) {
-		char line[40];
-
-		// Global or entry specific
-		if(i->entryIndex == 0)
-			sprintf(line, ".const %s constant0[%d] = {", i->type, i->m_constMemory.size());
-		else
-			sprintf(line, ".const %s ce%dc%d[%d] = {", i->type, i->entryIndex, i->index, i->m_constMemory.size());
-	
-		output(line);
-
-		std::list<std::string>::iterator j;
-		int l=0;
-		for(j=i->m_constMemory.begin(); j!=i->m_constMemory.end(); ++j) {
-			if(j!=i->m_constMemory.begin())
-				output(", ");
-			if( (l++ % 4) == 0) output("\n          ");
-			output(j->c_str());
-		}
-		output("\n};\n\n");
-	}
-
-
-	for(std::list<constMemory2>::iterator i=m_constMemoryList2.begin(); i!=m_constMemoryList2.end(); ++i) {
-		char line[1024];
-
-		// Global or entry specific
-		sprintf(line, ".const %s constant1%s[%d] = {", i->type, i->kernel, i->m_constMemory.size());
-
-		output(line);
-
-		std::list<std::string>::iterator j;
-		int l=0;
-		for(j=i->m_constMemory.begin(); j!=i->m_constMemory.end(); ++j) {
-			if(j!=i->m_constMemory.begin())
-				output(", ");
-			if( (l++ % 4) == 0) output("\n          ");
-			output(j->c_str());
-		}
-		output("\n};\n\n");
-	}
-
-	// Next, print out the local memory declaration
-	std::list<decudaEntry>::iterator e;
-	int eIndex=1; // entry index starts from 1 from the first blank entry is missing here (only in header entry list)
-	for(e=m_entryList.begin(); e!=m_entryList.end(); ++e) {
-		if(e->m_lMemSize > 0) {
-			std::stringstream out;
-			out << ".local .b8 l" << eIndex << "[" << e->m_lMemSize << "];" << std::endl;
-			output(out.str().c_str());
-		}
-		eIndex++;
-	}
-	output("\n");
-
-	// Next, print out the global memory declaration
-	std::list<globalMemory>::iterator g;
-	for(g=m_globalMemoryList.begin(); g!=m_globalMemoryList.end(); ++g) {
-		std::stringstream out;
-		out << ".global .b8 " << g->name << "[" << g->bytes << "];" << std::endl;
-		output(out.str().c_str());
-	}
-	output("\n");
-
-	// Next, print out constant memory pointers
-	std::list<constMemoryPtr>::iterator cp;
-	for(cp=m_constMemoryPtrList.begin(); cp!=m_constMemoryPtrList.end(); ++cp) {
-		std::stringstream out;
-		out << ".const .b8 " << cp->name << "[" << cp->bytes << "];" << std::endl;
-		out << ".constptr " << cp->name << ", " << cp->destination << ", " << cp->offset << ";" << std::endl;
-		output(out.str().c_str());
-	}
-	output("\n");
-
-}
-
-// add vector operand
-void decudaInstList::addVector(char* vector, int vectorSize) {
-	// If vector size is 1, make it 4 by adding blanks
-	if(vectorSize == 1) {
-		std::string vectorNew = vector;
-		vectorNew = vectorNew.substr(0,vectorNew.size()-1) + ",_,_,_}";
-		char* vectorNewName = new char [strlen(vectorNew.c_str())+1];
-		strcpy(vectorNewName, vectorNew.c_str());
-		getListEnd().addOperand(vectorNewName);
-	} else {
-		getListEnd().addOperand(vector);
-	}
-}
-
-
-// add memory operand
-// memType: 0=constant, 1=shared, 2=global, 3=local
-void decudaInstList::addMemoryOperand(std::string mem, int memType) {
-	std::string origMem = mem;
-	assert(0 && "This function is obsolete and shouldn't be used");
-
-	// If constant memory type, add prefix for entry specific constant memory
-	if(memType == 0) {
-		// Entry-specific constant memory c1
-		if(mem.substr(0, 3) == "c1[") {
-			std::stringstream out;
-			//out << "ce" << m_entryList.size() << mem;
-			mem = mem.substr(2);
-			out << "constant1" << this->m_entryList.back().m_entryName << mem;
-			mem = out.str();
-		}
-		// Global memory c14
-		// Replace this with the actual global memory name
-		else if(mem.substr(0, 3) == "c14") {
-			// Find the global memory identifier based on the offset provided
-			int offset;
-			sscanf(mem.substr(4,mem.size()-5).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) {
-				output("Could not find a global memory with this offset.\n");
-				assert(0);
-			}
-		}
-		// Global constant memory c0
-		else if(mem.substr(0, 3) == "c0[") {
-			// Do nothing
-		}
-		else {
-			output("Unrecognized memory type.\n");
-			assert(0);
-		}
-	}
-
-	// If local memory type, fix the decuda bug where l[4] is actually outputted l[$r4]
-	if(memType == 3) {
-		// Remove "$r" from "l[$r#]"
-		if(mem.substr(2, 2) == "$r") {
-			std::stringstream out;
-			out << mem.substr(0, 2) << mem.substr(4, mem.size()-4);
-			mem = out.str();
-		}
-
-		// Add entry entry number after 'l' to differentiate from other entries
-		std::stringstream out;
-		out << mem.substr(0,1) << m_entryList.size() << mem.substr(1,mem.size()-1);
-		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);
-}
-
-
-
-// get the list of real tex names
-std::list<std::string> decudaInstList::getRealTexList() {
-	return m_realTexList;
-}
-
-// set the list of real tex names
-void decudaInstList::setRealTexList(std::list<std::string> realTexList) {
-	m_realTexList = realTexList;
-}
-
-
-
-
-
-// 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 decudaInstList::readConstMemoryFromBinFile(std::string binString) {
-	// Initialize a list to store memory values
-	std::list<std::string> c0;
-	
-	// Get each constant segment
-	const boost::regex constPattern("(consts \\{[^\\}]*(mem \\{[^\\}]*\\})?[^\\}]*\\})");
-
-	// Parse each constseg
-	const boost::sregex_token_iterator end;
-	for (
-		boost::sregex_token_iterator i(binString.begin(),binString.end(), constPattern);
-		i != end;
-		++i
-	     )
-	{
-		// For each const segment, get the offset, bytes and memory values string
-		std::string constSeg_s = *i;
-		std::string offset_s, bytes_s, name, mem;
-		int offset, bytes;
-
-		boost::smatch offsetResult;
-		boost::smatch bytesResult;
-		boost::smatch nameResult;
-		boost::smatch memResult;
-
-		const boost::regex offsetPattern("offset\\s*=\\s(\\d*)");
-		const boost::regex bytesPattern("bytes\\s*=\\s(\\d*)");
-		const boost::regex namePattern("name\\s*=\\s(\\w*)");
-		const boost::regex memPattern("mem \\{([^\\}]*)\\}");
-
-		boost::regex_search(constSeg_s, offsetResult, offsetPattern);
-		boost::regex_search(constSeg_s, bytesResult, bytesPattern);
-		boost::regex_search(constSeg_s, nameResult, namePattern);
-		bool memExists = boost::regex_search(constSeg_s, memResult, memPattern);
-
-		//printf("\nmemexists=%d\n", memExists);
-
-		offset_s = offsetResult[1];
-		offset = atoi(offset_s.c_str());
-		bytes_s = bytesResult[1];
-		bytes = atoi(bytes_s.c_str());
-		name = nameResult[1];
-
-
-		// Resize the c0 list if needed
-		if(c0.size() < offset/4 + bytes/4) c0.resize(offset/4 + bytes/4, "0x00000000");
-
-		// If memory is initialized, import values
-		if(memExists) {
-			mem = memResult[1];
-
-			// Parse mem string, loop through each memory value and store it in the appropriate offset
-			// in the c0 list
-			// Before adding to the list, we increase the size of the list by inserting
-			// dummy elements. Then when adding memory values, the dummy elements are removed.
-			const boost::regex memValuePattern("(0x[A-Fa-f0-9]{8,8})");
-
-			// Initialize iterator
-			std::list<std::string>::iterator it = c0.begin();
-			std::advance(it, offset/4);
-
-			// Add values to memory list
-			const boost::sregex_token_iterator end2;
-			for (
-				boost::sregex_token_iterator j(mem.begin(),mem.end(), memValuePattern);
-				j != end2;
-				++j
-			     )
-			{
-				it = c0.erase(it);
-				c0.insert(it, *j);
-			}
-		} else {
-			// Uninitialized const memory - defined a const memory pointer
-			constMemoryPtr cMemPtr;
-			cMemPtr.bytes = bytes;
-			cMemPtr.offset = offset;
-			cMemPtr.name = name;
-			cMemPtr.destination = "c0";
-
-			m_constMemoryPtrList.push_back(cMemPtr);
-		}
-	}
-
-	
-	// Finished parsing of the file, now iterate over the list and add values to constant memory segment
-	if(c0.size() > 0) {
-		addConstMemory(0);
-		setConstMemoryType(".u32");	
-		std::list<std::string>::iterator c;
-		for(c=c0.begin(); c!=c0.end(); ++c) {
-			addConstMemoryValue(*c);
-		}
-	}		
-}
-
-// Read in global memory from bin file
-void decudaInstList::readGlobalMemoryFromBinFile(std::string binString) {
-	// Get each constant segment
-	const boost::regex globalPattern("(reloc \\{[^\\}]*segnum  = 14[^\\}]*\\})");
-
-	// Parse each constseg
-	const boost::sregex_token_iterator end;
-	for (
-		boost::sregex_token_iterator i(binString.begin(),binString.end(), globalPattern);
-		i != end;
-		++i
-	     )
-	{
-		// For each global segment, get the offset, bytes and name
-		std::string globalSeg_s = *i;
-		std::string offset_s, bytes_s, name;
-		int offset, bytes;
-
-		boost::smatch offsetResult;
-		boost::smatch bytesResult;
-		boost::smatch nameResult;
-
-		const boost::regex offsetPattern("offset\\s*=\\s(\\d*)");
-		const boost::regex bytesPattern("bytes\\s*=\\s(\\d*)");
-		const boost::regex namePattern("name\\s*=\\s(\\w*)");
-
-		boost::regex_search(globalSeg_s, offsetResult, offsetPattern);
-		boost::regex_search(globalSeg_s, bytesResult, bytesPattern);
-		boost::regex_search(globalSeg_s, nameResult, namePattern);
-
-		offset_s = offsetResult[1];
-		offset = atoi(offset_s.c_str());
-		bytes_s = bytesResult[1];
-		bytes = atoi(bytes_s.c_str());
-		name = nameResult[1];
-
-		// Add global memory
-		globalMemory gMem;
-		gMem.offset = offset;
-		gMem.bytes = bytes;
-		gMem.name = name;
-
-		m_globalMemoryList.push_back(gMem);
-	}
-}
diff --git a/decuda_to_ptxplus/decudaInstList.h b/decuda_to_ptxplus/decudaInstList.h
deleted file mode 100644
index d7387ee..0000000
--- a/decuda_to_ptxplus/decudaInstList.h
+++ /dev/null
@@ -1,187 +0,0 @@
-// Copyright (c) 2009-2011, Jimmy Kwa,
-// The University of British Columbia
-// All rights reserved.
-//
-// Redistribution and use in source and binary forms, with or without
-// modification, are permitted provided that the following conditions are met:
-//
-// Redistributions of source code must retain the above copyright notice, this
-// list of conditions and the following disclaimer.
-// Redistributions in binary form must reproduce the above copyright notice, this
-// list of conditions and the following disclaimer in the documentation and/or
-// other materials provided with the distribution.
-// Neither the name of The University of British Columbia nor the names of its
-// contributors may be used to endorse or promote products derived from this
-// software without specific prior written permission.
-//
-// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
-// ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
-// WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
-// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
-// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
-// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
-// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
-// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
-// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
-// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-
-#include "decudaInst.h"
-#include <list>
-#include <string>
-#include <cstring>
-#include <stdio.h>
-#include <iostream>
-#include <sstream>
-#include <stdlib.h>
-#include <boost/regex.hpp>
-#include <map>
-
-// Used for entry specific constant memory segments (c1)
-struct constMemory
-{
-	int index;
-	int entryIndex;
-	const char* type;
-	std::list<std::string> m_constMemory;
-};
-
-struct constMemory2
-{
-	const char* kernel;
-	const char* type;
-	std::list<std::string> m_constMemory;
-};
-
-// Used for uninitialized constant memory (globally defined)
-struct constMemoryPtr
-{
-	int bytes;
-	std::string name;
-
-	std::string destination;
-	int offset;
-};
-
-// Used for global memory segments
-struct globalMemory
-{
-	int offset;
-	int bytes;
-	std::string name;
-};
-
-
-
-struct decudaEntry
-{
-	//char* m_entryName;
-	std::string m_entryName;
-	std::list<decudaInst> m_instList;	// List of decuda instructions
-
-	// Register list
-	int m_largestRegIndex;
-	int m_largestOfsRegIndex;
-	bool m_reg124;
-	bool m_oreg127;
-
-	// Predicate list
-	int m_largestPredIndex;
-
-	// Local memory size
-	int m_lMemSize;
-
-	//use for recording used labels
-	std::list<std::string> m_labelList;
-
-   // Histogram for operation per cycle count
-   std::map<std::string, int> m_opPerCycleHistogram; 
-};
-
-
-
-class decudaInstList
-{
-
-protected:
-	// List of decuda entries
-	std::list<decudaEntry> m_entryList;
-
-	// Const memory list
-	std::list<constMemory> m_constMemoryList;
-	std::list<constMemory2> m_constMemoryList2;
-
-	// Const memory pointers list
-	std::list<constMemoryPtr> m_constMemoryPtrList;
-
-	// Global memory list
-	std::list<globalMemory> m_globalMemoryList;
-
-	// Tex list
-	std::list<std::string>  m_realTexList;	// Stores the real names of tex variables
-
-	// Print register names
-	void printRegNames(decudaEntry entry);
-	void printOutOfBoundRegisters(decudaEntry entry);
-
-	// Print predicate names
-	void printPredNames(decudaEntry entry);
-
-	// Print const memory directives
-	void printMemory();
-
-	// Increment register or predicate offsets
-	std::string parseRegister(std::string reg, bool lo=false, int vectorFlag=0);
-	std::string parsePredicate(std::string pred);
-
-public:
-	//constructor
-	decudaInstList();
-
-	//accessors
-	decudaInst getListEnd();
-
-	//mutator
-	int addEntry(std::string entryName); // creates a new entry
-	void setLastEntryName(std::string entryName); // sets name of last entry
-	void setLastEntryLMemSize(int lMemSize); // sets the local memory size of last entry
-	bool findEntry(std::string entryName, decudaEntry& entry); // find and return entry
-	
-
-	int add(decudaInst* newInst); //add DecudaInst to list
-
-	void addRegister(std::string reg, bool lo=false); //add register
-	void addPredicate(std::string pred); //add predicate
-	void addDoublePredReg(std::string pred, std::string reg, bool lo=false); // add pred|reg double operand
-
-	void addTex(std::string tex);	// add tex operand
-
-	void addVector(char* vector, int vectorSize); // add vector operand
-
-	void addMemoryOperand(std::string mem, int memType); // add memory operand
-
-
-	// Parsing constant memory segments list
-	void addEntryConstMemory(int index); // add entry specific const memory
-	void addConstMemory(int index); // add global const memory
-	void setConstMemoryType(const char* type); // set type of constant memory
-	void setConstMemoryType2(const char* type); // set type of constant memory
-	void addConstMemoryValue(std::string constMemoryValue); // add const memory
-	void addConstMemoryValue2(std::string constMemoryValue); // add const memory
-
-	std::list<std::string> getRealTexList(); // get the list of real tex names
-	void setRealTexList(std::list<std::string> realTexList); // set the list of real tex names
-
-	void readConstMemoryFromBinFile(std::string binString); // read in constant memory from bin file
-	void readGlobalMemoryFromBinFile(std::string binString); // read in global memory from bin file
-
-	//print representation
-	void printHeaderInstList();
-	void printDecudaInstList();
-	void printNewPtxList(decudaInstList* headerInfo);
-
-
-	// debug helper methods
-	void printEntryNames();
-
-};
diff --git a/decuda_to_ptxplus/decuda_to_ptxplus.cc b/decuda_to_ptxplus/decuda_to_ptxplus.cc
deleted file mode 100644
index ad48963..0000000
--- a/decuda_to_ptxplus/decuda_to_ptxplus.cc
+++ /dev/null
@@ -1,115 +0,0 @@
-// Copyright (c) 2009-2011, Jimmy Kwa,
-// The University of British Columbia
-// All rights reserved.
-//
-// Redistribution and use in source and binary forms, with or without
-// modification, are permitted provided that the following conditions are met:
-//
-// Redistributions of source code must retain the above copyright notice, this
-// list of conditions and the following disclaimer.
-// Redistributions in binary form must reproduce the above copyright notice, this
-// list of conditions and the following disclaimer in the documentation and/or
-// other materials provided with the distribution.
-// Neither the name of The University of British Columbia nor the names of its
-// contributors may be used to endorse or promote products derived from this
-// software without specific prior written permission.
-//
-// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
-// ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
-// WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
-// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
-// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
-// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
-// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
-// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
-// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
-// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-
-#include <iostream>
-#include "decudaInstList.h"
-#include <stdio.h>
-#include<fstream>
-
-using namespace std;
-
-decudaInstList *g_instList = new decudaInstList();
-decudaInstList *g_headerList = new decudaInstList();
-
-int yyparse();
-extern "C" FILE *yyin;
-
-int ptx_parse();
-extern "C" FILE *ptx_in;
-
-FILE *bin_in;
-FILE *ptxplus_out;
-
-void output(const char * 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;
-}
-
-int main(int argc, char* argv[])
-{
-	if(argc != 5)
-	{
-		cout << "Usage: decuda_to_ptxplus [decuda filename] [ptx filename] [bin filename] [ptxplus output filename]\n";
-		return 0;
-	}
-
-	const char *decudaFilename = argv[1];
-	const char *ptxFilename = argv[2];
-	const char *binFilename = argv[3];
-	const char *ptxplusFilename = argv[4];
-
-	//header_in = fopen( ptxFilename, "r" );
-	ptx_in = fopen( ptxFilename, "r" );
-	yyin = fopen( decudaFilename, "r" );
-	bin_in = fopen( binFilename, "r" );
-	ptxplus_out = fopen( ptxplusFilename, "w" );
-
-
-	fileToString(binFilename);
-
-	printf("RUNNING decuda2ptxplus ...\n");
-
-	// Parse original ptx
-	ptx_parse();
-
-	// Copy real tex list from ptx to ptxplus instruction list
-	g_instList->setRealTexList(g_headerList->getRealTexList());
-
-	// Insert constant memory from bin file
-	g_instList->readConstMemoryFromBinFile(fileToString(binFilename));
-
-	// Insert global memory from bin file
-	g_instList->readGlobalMemoryFromBinFile(fileToString(binFilename));
-
-	// Parse decuda output
-	yyparse();
-	printf("END RUN\n");
-
-	// Print ptxplus
-	g_headerList->printHeaderInstList();
-	g_instList->printNewPtxList(g_headerList);
-
-	fclose(ptx_in);
-	fclose(yyin);
-	fclose(bin_in);
-	fclose(ptxplus_out);
-
-	printf("DONE. \n");
-
-	return 0;
-}
diff --git a/decuda_to_ptxplus/makefile b/decuda_to_ptxplus/makefile
deleted file mode 100644
index ea4867f..0000000
--- a/decuda_to_ptxplus/makefile
+++ /dev/null
@@ -1,159 +0,0 @@
-# Copyright (c) 2009-2011, Jimmy Kwa,
-# The University of British Columbia
-# All rights reserved.
-#
-# Redistribution and use in source and binary forms, with or without
-# modification, are permitted provided that the following conditions are met:
-#
-# Redistributions of source code must retain the above copyright notice, this
-# list of conditions and the following disclaimer.
-# Redistributions in binary form must reproduce the above copyright notice, this
-# list of conditions and the following disclaimer in the documentation and/or
-# other materials provided with the distribution.
-# Neither the name of The University of British Columbia nor the names of its
-# contributors may be used to endorse or promote products derived from this
-# software without specific prior written permission.
-#
-# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
-# ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
-# WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
-# DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
-# FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
-# DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
-# SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
-# CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
-# OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
-# OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-
-FLAGS = -ggdb -I${BOOST_ROOT} -lboost_regex -lboost_filesystem -L${BOOST_LIB}
-
-all: decuda_to_ptxplus
-
-decuda_to_ptxplus: stringList.o decudaInst.o decudaInstList.o decudaLexer.o decudaParser.o headerParser.o headerLexer.o decuda_to_ptxplus.o lex.ptx_.o ptx.tab.o
-	g++ ${FLAGS} -odecuda_to_ptxplus stringList.o decudaInst.o decudaInstList.o decudaLexer.o decudaParser.o headerParser.o headerLexer.o decuda_to_ptxplus.o lex.ptx_.o ptx.tab.o
-
-stringList.o : stringList.cc stringList.h
-	g++ ${FLAGS} -c -ostringList.o stringList.cc
-
-decudaInst.o :  decudaInst.cc decudaInst.h stringList.h
-	g++ ${FLAGS} -c -odecudaInst.o decudaInst.cc
-
-decudaInstList.o : decudaInstList.cc decudaInstList.h decudaInst.h
-	g++ ${FLAGS} -c -odecudaInstList.o decudaInstList.cc
-
-decudaParser.cpp : decuda.y decudaInstList.h
-	bison -d -odecudaParser.cpp decuda.y
-
-decudaParser.hpp : decudaParser.cpp
-
-decudaLexer.cpp : decuda.l decudaParser.hpp
-	flex -B -odecudaLexer.cpp decuda.l
-
-headerParser.cpp : header.y decudaInstList.h
-	bison --name-prefix=header_ -d -oheaderParser.cpp header.y
-
-headerParser.hpp : headerParser.cpp
-
-headerLexer.cpp : header.l headerParser.hpp
-	flex -B -oheaderLexer.cpp header.l
-
-decudaLexer.o : decudaLexer.cpp
-	g++ ${FLAGS} -c -Igenerated -odecudaLexer.o decudaLexer.cpp
-
-decudaParser.o : decudaParser.cpp stringList.h decudaInst.h decudaInstList.h 
-	g++ ${FLAGS} -c -Igenerated -odecudaParser.o decudaParser.cpp
-
-headerLexer.o : headerLexer.cpp
-	g++ ${FLAGS} -c -Igenerated -oheaderLexer.o headerLexer.cpp
-
-headerParser.o : headerParser.cpp stringList.h decudaInst.h decudaInstList.h 
-	g++ ${FLAGS} -c -Igenerated -oheaderParser.o headerParser.cpp
-
-decuda_to_ptxplus.o : decuda_to_ptxplus.cc decudaInstList.h
-	g++ ${FLAGS} -c -odecuda_to_ptxplus.o decuda_to_ptxplus.cc
-
-lex.ptx_.c : ../src/cuda-sim/ptx.l
-	flex ../src/cuda-sim/ptx.l
-
-ptx.tab.c : ../src/cuda-sim/ptx.y
-	bison --name-prefix=ptx_ -v -d ../src/cuda-sim/ptx.y
-
-ptx.tab.h : ptx.tab.c
-
-lex.ptx_.o : lex.ptx_.c ptx.tab.h
-	g++ ${FLAGS} -c lex.ptx_.c
-
-ptx.tab.o : ptx.tab.c
-	g++ ${FLAGS} -c ptx.tab.c
-
-SRCS = $(shell ls *.cc)
-
-depend:
-	makedepend $(SRCS) 2> /dev/null
-	
-clean:
-	rm -f ./stringList.o ./decudaInst.o ./decudaInstList.o ./decudaParser.o ./decudaLexer.o ./headerParser.o ./headerLexer.o ./decuda_to_ptxplus.o
-	rm -f ./decudaParser.hpp ./headerParser.hpp
-	rm -f ./decudaParser.cpp ./decudaLexer.cpp ./headerParser.cpp ./headerLexer.cpp
-	rm -f ./decuda_to_ptxplus
-	rm -f ./ptx.output ./ptx.tab.c ./ptx.tab.h ./ptx.tab.o
-	rm -f ./lex.ptx_.c ./lex.ptx_.o
-# DO NOT DELETE
-
-decudaInst.o: decudaInst.h stringList.h /usr/include/assert.h
-decudaInst.o: /usr/include/features.h /usr/include/sys/cdefs.h
-decudaInst.o: /usr/include/bits/wordsize.h /usr/include/gnu/stubs.h
-decudaInst.o: /usr/include/gnu/stubs-64.h /usr/include/stdlib.h
-decudaInst.o: /usr/include/bits/waitflags.h /usr/include/bits/waitstatus.h
-decudaInst.o: /usr/include/endian.h /usr/include/bits/endian.h
-decudaInst.o: /usr/include/bits/byteswap.h /usr/include/sys/types.h
-decudaInst.o: /usr/include/bits/types.h /usr/include/bits/typesizes.h
-decudaInst.o: /usr/include/time.h /usr/include/sys/select.h
-decudaInst.o: /usr/include/bits/select.h /usr/include/bits/sigset.h
-decudaInst.o: /usr/include/bits/time.h /usr/include/sys/sysmacros.h
-decudaInst.o: /usr/include/bits/pthreadtypes.h /usr/include/alloca.h
-decudaInst.o: /usr/include/stdio.h /usr/include/libio.h
-decudaInst.o: /usr/include/_G_config.h /usr/include/wchar.h
-decudaInst.o: /usr/include/bits/stdio_lim.h /usr/include/bits/sys_errlist.h
-decudaInstList.o: decudaInstList.h decudaInst.h stringList.h
-decudaInstList.o: /usr/include/assert.h /usr/include/features.h
-decudaInstList.o: /usr/include/sys/cdefs.h /usr/include/bits/wordsize.h
-decudaInstList.o: /usr/include/gnu/stubs.h /usr/include/gnu/stubs-64.h
-decudaInstList.o: /usr/include/stdio.h /usr/include/bits/types.h
-decudaInstList.o: /usr/include/bits/typesizes.h /usr/include/libio.h
-decudaInstList.o: /usr/include/_G_config.h /usr/include/wchar.h
-decudaInstList.o: /usr/include/bits/stdio_lim.h
-decudaInstList.o: /usr/include/bits/sys_errlist.h /usr/include/stdlib.h
-decudaInstList.o: /usr/include/bits/waitflags.h
-decudaInstList.o: /usr/include/bits/waitstatus.h /usr/include/endian.h
-decudaInstList.o: /usr/include/bits/endian.h /usr/include/bits/byteswap.h
-decudaInstList.o: /usr/include/sys/types.h /usr/include/time.h
-decudaInstList.o: /usr/include/sys/select.h /usr/include/bits/select.h
-decudaInstList.o: /usr/include/bits/sigset.h /usr/include/bits/time.h
-decudaInstList.o: /usr/include/sys/sysmacros.h
-decudaInstList.o: /usr/include/bits/pthreadtypes.h /usr/include/alloca.h
-decuda_to_ptxplus.o: decudaInstList.h decudaInst.h stringList.h
-decuda_to_ptxplus.o: /usr/include/assert.h /usr/include/features.h
-decuda_to_ptxplus.o: /usr/include/sys/cdefs.h /usr/include/bits/wordsize.h
-decuda_to_ptxplus.o: /usr/include/gnu/stubs.h /usr/include/gnu/stubs-64.h
-decuda_to_ptxplus.o: /usr/include/stdio.h /usr/include/bits/types.h
-decuda_to_ptxplus.o: /usr/include/bits/typesizes.h /usr/include/libio.h
-decuda_to_ptxplus.o: /usr/include/_G_config.h /usr/include/wchar.h
-decuda_to_ptxplus.o: /usr/include/bits/stdio_lim.h
-decuda_to_ptxplus.o: /usr/include/bits/sys_errlist.h /usr/include/stdlib.h
-decuda_to_ptxplus.o: /usr/include/bits/waitflags.h
-decuda_to_ptxplus.o: /usr/include/bits/waitstatus.h /usr/include/endian.h
-decuda_to_ptxplus.o: /usr/include/bits/endian.h /usr/include/bits/byteswap.h
-decuda_to_ptxplus.o: /usr/include/sys/types.h /usr/include/time.h
-decuda_to_ptxplus.o: /usr/include/sys/select.h /usr/include/bits/select.h
-decuda_to_ptxplus.o: /usr/include/bits/sigset.h /usr/include/bits/time.h
-decuda_to_ptxplus.o: /usr/include/sys/sysmacros.h
-decuda_to_ptxplus.o: /usr/include/bits/pthreadtypes.h /usr/include/alloca.h
-stringList.o: stringList.h /usr/include/stdio.h /usr/include/features.h
-stringList.o: /usr/include/sys/cdefs.h /usr/include/bits/wordsize.h
-stringList.o: /usr/include/gnu/stubs.h /usr/include/gnu/stubs-64.h
-stringList.o: /usr/include/bits/types.h /usr/include/bits/typesizes.h
-stringList.o: /usr/include/libio.h /usr/include/_G_config.h
-stringList.o: /usr/include/wchar.h /usr/include/bits/stdio_lim.h
-stringList.o: /usr/include/bits/sys_errlist.h
diff --git a/decuda_to_ptxplus/opcodes.def b/decuda_to_ptxplus/opcodes.def
deleted file mode 100644
index 7815f24..0000000
--- a/decuda_to_ptxplus/opcodes.def
+++ /dev/null
@@ -1,159 +0,0 @@
-/* 
- * Copyright (c) 2009 by Tor M. Aamodt, Wilson W. L. Fung, Ali Bakhoda, 
- * George L. Yuan and the University of British Columbia
- * Vancouver, BC  V6T 1Z4
- * All Rights Reserved.
- * 
- * THIS IS A LEGAL DOCUMENT BY DOWNLOADING GPGPU-SIM, YOU ARE AGREEING TO THESE
- * TERMS AND CONDITIONS.
- * 
- * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
- * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
- * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
- * ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNERS OR CONTRIBUTORS BE
- * LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
- * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
- * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
- * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
- * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
- * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
- * POSSIBILITY OF SUCH DAMAGE.
- * 
- * NOTE: The files libcuda/cuda_runtime_api.c and src/cuda-sim/cuda-math.h
- * are derived from the CUDA Toolset available from http://www.nvidia.com/cuda
- * (property of NVIDIA).  The files benchmarks/BlackScholes/ and 
- * benchmarks/template/ are derived from the CUDA SDK available from 
- * http://www.nvidia.com/cuda (also property of NVIDIA).  The files from 
- * src/intersim/ are derived from Booksim (a simulator provided with the 
- * textbook "Principles and Practices of Interconnection Networks" available 
- * from http://cva.stanford.edu/books/ppin/). As such, those files are bound by 
- * the corresponding legal terms and conditions set forth separately (original 
- * copyright notices are left in files from these sources and where we have 
- * modified a file our copyright notice appears before the original copyright 
- * notice).  
- * 
- * Using this version of GPGPU-Sim requires a complete installation of CUDA 
- * which is distributed seperately by NVIDIA under separate terms and 
- * conditions.  To use this version of GPGPU-Sim with OpenCL requires a
- * recent version of NVIDIA's drivers which support OpenCL.
- * 
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions are met:
- * 
- * 1. Redistributions of source code must retain the above copyright notice,
- * this list of conditions and the following disclaimer.
- * 
- * 2. Redistributions in binary form must reproduce the above copyright notice,
- * this list of conditions and the following disclaimer in the documentation
- * and/or other materials provided with the distribution.
- * 
- * 3. Neither the name of the University of British Columbia nor the names of
- * its contributors may be used to endorse or promote products derived from
- * this software without specific prior written permission.
- * 
- * 4. This version of GPGPU-SIM is distributed freely for non-commercial use only.  
- *  
- * 5. No nonprofit user may place any restrictions on the use of this software,
- * including as modified by the user, by any other authorized user.
- * 
- * 6. GPGPU-SIM was developed primarily by Tor M. Aamodt, Wilson W. L. Fung, 
- * Ali Bakhoda, George L. Yuan, at the University of British Columbia, 
- * Vancouver, BC V6T 1Z4
- */
-
-/*6th operand of each OP_DEF  reflects its classification  */
-/*Type 
-ALU 1
-MAD 2
-Control 3
-SFU 4
-Mem(except Tex) 5
-Tex 6
-Nop 7
-Other 10
-*/
-OP_DEF(ABS_OP,abs_impl,"abs",1,1)
-OP_DEF(ADD_OP,add_impl,"add",1,1)
-OP_DEF(ADDC_OP,addc_impl,"addc",1,1)
-OP_DEF(AND_OP,and_impl,"and",1,1)
-OP_DEF(ANDN_OP,andn_impl,"andn",1,1)
-OP_DEF(ATOM_OP,atom_impl,"atom",0,3)
-OP_DEF(BAR_OP,bar_sync_impl,"bar.sync",1,3)
-OP_DEF(BFE_OP,bfe_impl,"bfe",1,1)
-OP_DEF(BFI_OP,bfi_impl,"bfi",1,1)
-OP_DEF(BFIND_OP,bfind_impl,"bfind",1,1)
-OP_DEF(BRA_OP,bra_impl,"bra",0,3)
-OP_DEF(BREV_OP,brev_impl,"brev",1,1)
-OP_DEF(BRKPT_OP,brkpt_impl,"brkpt",1,9)
-OP_DEF(CALL_OP,call_impl,"call",1,3)
-OP_DEF(CALLP_OP,callp_impl,"callp",1,3)
-OP_DEF(CLZ_OP,clz_impl,"clz",1,1)
-OP_DEF(CNOT_OP,cnot_impl,"cnot",1,1)
-OP_DEF(COS_OP,cos_impl,"cos",1,4)
-OP_DEF(CVT_OP,cvt_impl,"cvt",1,1)
-OP_DEF(CVTA_OP,cvta_impl,"cvta",1,1)
-OP_DEF(DIV_OP,div_impl,"div",1,1)
-OP_DEF(EX2_OP,ex2_impl,"ex2",1,4)
-OP_DEF(EXIT_OP,exit_impl,"exit",1,3)
-OP_DEF(FMA_OP,fma_impl,"fma",1,2)
-OP_DEF(ISSPACEP_OP,isspacep_impl,"isspacep",1,1)
-OP_DEF(LD_OP,ld_impl,"ld",1,5)
-OP_DEF(LDU_OP,ldu_impl,"ldu",1,5)
-OP_DEF(LG2_OP,lg2_impl,"lg2",1,4)
-OP_DEF(MAD24_OP,mad24_impl,"mad24",1,2)
-OP_DEF(MAD_OP,mad_impl,"mad",1,2)
-OP_DEF(MAX_OP,max_impl,"max",1,1)
-OP_DEF(MEMBAR_OP,membar_impl,"membar",1,3)
-OP_DEF(MIN_OP,min_impl,"min",1,1)
-OP_DEF(MOV_OP,mov_impl,"mov",1,1)
-OP_DEF(MUL24_OP,mul24_impl,"mul24",1,1)
-OP_DEF(MUL_OP,mul_impl,"mul",1,1)
-OP_DEF(NEG_OP,neg_impl,"neg",1,1)
-OP_DEF(NANDN_OP,nandn_impl,"nandn",1,1)
-OP_DEF(NORN_OP,norn_impl,"norn",1,1)
-OP_DEF(NOT_OP,not_impl,"not",1,1)
-OP_DEF(OR_OP,or_impl,"or",1,1)
-OP_DEF(ORN_OP,orn_impl,"orn",1,1)
-OP_DEF(PMEVENT_OP,pmevent_impl,"pmevent",1,10)
-OP_DEF(POPC_OP,popc_impl,"popc",1,1)
-OP_DEF(PREFETCH_OP,prefetch_impl,"prefetch",1,5)
-OP_DEF(PREFETCHU_OP,prefetchu_impl,"prefetchu",1,5)
-OP_DEF(PRMT_OP,prmt_impl,"prmt",1,1)
-OP_DEF(RCP_OP,rcp_impl,"rcp",1,4)
-OP_DEF(RED_OP,red_impl,"red",1,7)
-OP_DEF(REM_OP,rem_impl,"rem",1,1)
-OP_DEF(RET_OP,ret_impl,"ret",0,3)
-OP_DEF(RETP_OP,retp_impl,"retp",0,3)
-OP_DEF(RSQRT_OP,rsqrt_impl,"rsqrt",1,4)
-OP_DEF(SAD_OP,sad_impl,"sad",1,1)
-OP_DEF(SELP_OP,selp_impl,"selp",1,1)
-OP_DEF(SETP_OP,setp_impl,"setp",1,1)
-OP_DEF(SET_OP,set_impl,"set",1,1)
-OP_DEF(SHL_OP,shl_impl,"shl",1,1)
-OP_DEF(SHR_OP,shr_impl,"shr",1,1)
-OP_DEF(SIN_OP,sin_impl,"sin",1,4)
-OP_DEF(SLCT_OP,slct_impl,"slct",1,1)
-OP_DEF(SQRT_OP,sqrt_impl,"sqrt",1,4)
-OP_DEF(ST_OP,st_impl,"st",0,5)
-OP_DEF(SUB_OP,sub_impl,"sub",1,1)
-OP_DEF(SUBC_OP,subc_impl,"subc",1,1)
-OP_DEF(SULD_OP,suld_impl,"suld",1,6)
-OP_DEF(SURED_OP,sured_impl,"sured",1,6)
-OP_DEF(SUST_OP,sust_impl,"sust",1,6)
-OP_DEF(SUQ_OP,suq_impl,"suq",1,6)
-OP_DEF(TEX_OP,tex_impl,"tex",1,6)
-OP_DEF(TRAP_OP,trap_impl,"trap",1,3)
-OP_DEF(VABSDIFF_OP,vabsdiff_impl,"vabsdiff",0,11)
-OP_DEF(VADD_OP,vadd_impl,"vadd",0,11)
-OP_DEF(VMAD_OP,vmad_impl,"vmad",0,11)
-OP_DEF(VMAX_OP,vmax_impl,"vmax",0,11)
-OP_DEF(VMIN_OP,vmin_impl,"vmin",0,11)
-OP_DEF(VSET_OP,vset_impl,"vset",0,11)
-OP_DEF(VSHL_OP,vshl_impl,"vshl",0,11)
-OP_DEF(VSHR_OP,vshr_impl,"vshr",0,11)
-OP_DEF(VSUB_OP,vsub_impl,"vsub",0,11)
-OP_DEF(VOTE_OP,vote_impl,"vote",0,3)
-OP_DEF(XOR_OP,xor_impl,"xor",1,1)
-OP_DEF(NOP_OP,nop_impl,"nop",0,7)
-OP_DEF(BREAK_OP,break_impl,"break",0,3)
-OP_DEF(BREAKADDR_OP,breakaddr_impl,"breakaddr",0,3)
diff --git a/decuda_to_ptxplus/opcodes.h b/decuda_to_ptxplus/opcodes.h
deleted file mode 100644
index 2ceafaf..0000000
--- a/decuda_to_ptxplus/opcodes.h
+++ /dev/null
@@ -1,97 +0,0 @@
-/* 
- * Copyright (c) 2009 by Tor M. Aamodt and the University of British Columbia
- * Vancouver, BC  V6T 1Z4
- * All Rights Reserved.
- * 
- * THIS IS A LEGAL DOCUMENT BY DOWNLOADING GPGPU-SIM, YOU ARE AGREEING TO THESE
- * TERMS AND CONDITIONS.
- * 
- * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
- * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
- * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
- * ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNERS OR CONTRIBUTORS BE
- * LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
- * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
- * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
- * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
- * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
- * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
- * POSSIBILITY OF SUCH DAMAGE.
- * 
- * NOTE: The files libcuda/cuda_runtime_api.c and src/cuda-sim/cuda-math.h
- * are derived from the CUDA Toolset available from http://www.nvidia.com/cuda
- * (property of NVIDIA).  The files benchmarks/BlackScholes/ and 
- * benchmarks/template/ are derived from the CUDA SDK available from 
- * http://www.nvidia.com/cuda (also property of NVIDIA).  The files from 
- * src/intersim/ are derived from Booksim (a simulator provided with the 
- * textbook "Principles and Practices of Interconnection Networks" available 
- * from http://cva.stanford.edu/books/ppin/). As such, those files are bound by 
- * the corresponding legal terms and conditions set forth separately (original 
- * copyright notices are left in files from these sources and where we have 
- * modified a file our copyright notice appears before the original copyright 
- * notice).  
- * 
- * Using this version of GPGPU-Sim requires a complete installation of CUDA 
- * which is distributed seperately by NVIDIA under separate terms and 
- * conditions.  To use this version of GPGPU-Sim with OpenCL requires a
- * recent version of NVIDIA's drivers which support OpenCL.
- * 
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions are met:
- * 
- * 1. Redistributions of source code must retain the above copyright notice,
- * this list of conditions and the following disclaimer.
- * 
- * 2. Redistributions in binary form must reproduce the above copyright notice,
- * this list of conditions and the following disclaimer in the documentation
- * and/or other materials provided with the distribution.
- * 
- * 3. Neither the name of the University of British Columbia nor the names of
- * its contributors may be used to endorse or promote products derived from
- * this software without specific prior written permission.
- * 
- * 4. This version of GPGPU-SIM is distributed freely for non-commercial use only.  
- *  
- * 5. No nonprofit user may place any restrictions on the use of this software,
- * including as modified by the user, by any other authorized user.
- * 
- * 6. GPGPU-SIM was developed primarily by Tor M. Aamodt, Wilson W. L. Fung, 
- * Ali Bakhoda, George L. Yuan, at the University of British Columbia, 
- * Vancouver, BC V6T 1Z4
- */
-
-#ifndef opcodes_h_included
-#define opcodes_h_included
-
-enum opcode_t {
-#define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) OP,
-#include "opcodes.def"
-   NUM_OPCODES
-#undef OP_DEF
-};
-
-enum special_regs {
-   CLOCK_REG,
-   HALFCLOCK_ID,
-   CLOCK64_REG,
-   CTAID_REG,
-   ENVREG_REG,
-   GRIDID_REG,
-   LANEID_REG,
-   LANEMASK_EQ_REG,
-   LANEMASK_LE_REG,
-   LANEMASK_LT_REG,
-   LANEMASK_GE_REG,
-   LANEMASK_GT_REG,
-   NCTAID_REG,
-   NTID_REG,
-   NSMID_REG,
-   NWARPID_REG,
-   PM_REG,
-   SMID_REG,
-   TID_REG,
-   WARPID_REG,
-   WARPSZ_REG
-};
-
-#endif
diff --git a/decuda_to_ptxplus/ptx_parser.h b/decuda_to_ptxplus/ptx_parser.h
deleted file mode 100644
index bed82bb..0000000
--- a/decuda_to_ptxplus/ptx_parser.h
+++ /dev/null
@@ -1,329 +0,0 @@
-// Copyright (c) 2009-2011, Jimmy Kwa,
-// The University of British Columbia
-// All rights reserved.
-//
-// Redistribution and use in source and binary forms, with or without
-// modification, are permitted provided that the following conditions are met:
-//
-// Redistributions of source code must retain the above copyright notice, this
-// list of conditions and the following disclaimer.
-// Redistributions in binary form must reproduce the above copyright notice, this
-// list of conditions and the following disclaimer in the documentation and/or
-// other materials provided with the distribution.
-// Neither the name of The University of British Columbia nor the names of its
-// contributors may be used to endorse or promote products derived from this
-// software without specific prior written permission.
-//
-// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
-// ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
-// WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
-// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
-// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
-// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
-// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
-// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
-// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
-// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-
-   #include <cstdlib>
-   #include <cstring>
-   #include <string>
-   #include <list>
-   #include <map>
-   #include <vector>
-   #include <assert.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:%u) ", __FILE__, __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 ) {}
-void end_function() {}
-void add_directive() {}
-void add_function_arg() {}
-void add_instruction() {}
-void add_file( unsigned a, const char *b ) {}
-void add_variables() {}
-void set_variable_type() {}
-void add_identifier( const char *a, int b, unsigned c ) {}
-void add_option(int a ) {}
-void add_array_initializer() {}
-void add_label( const char *a ) {}
-void set_return() {}
-void add_opcode( int a ) {}
-void add_pred( const char *a, int b, int c ) {}
-void add_scalar_operand( const char *a ) {}
-void add_neg_pred_operand( const char *a ) {}
-void add_address_operand( const char *a, int b ) {}
-void change_operand_lohi( int a ) {}
-void change_double_operand_type( int a ) {}
-void change_operand_neg( ) {}
-void add_double_operand( const char *a, const char *b ) {}
-void add_2vector_operand( const char *a, const char *b ) {}
-void add_3vector_operand( const char *a, const char *b, const char *c ) {}
-void add_4vector_operand( const char *a, const char *b, const char *c, const char *d ) {}
-void add_builtin_operand( int a, int b ) {}
-void add_memory_operand() {}
-void change_memory_addr_space( const char *a ) {}
-void add_literal_int( int a ) {}
-void add_literal_float( float a ) {}
-void add_literal_double( double a ) {}
-void func_header_info_int(const char*, int) {}
-void add_extern_spec() {}
-void add_alignment_spec( int ) {}
-void add_pragma( const char *a ) {}
-
-void add_constptr(const char* identifier1, const char* identifier2, int offset) {}
-
-/*non-dummy stuff below this point*/
-
-#include <iostream>
-#include <sstream>
-#include "decudaInstList.h"
-#include <string.h>
-
-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;
-
-// Global variable to track if we are currently inside a tex directive
-bool inTexDirective = false;
-
-void add_function_name( const char *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("add_space_spec: spec=%u", spec);
-	decudaInst *instEntry;
-	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) {
-				instEntry = new decudaInst();
-				instEntry->setBase(".const");
-				g_headerList->add(instEntry);
-			}
-			break;
-	}
-}
-
-void add_scalar_type_spec( int headerInput )
-{
-	const char* compareString = g_headerList->getListEnd().getBase();
-
-	if( (inEntryDirective && inParamDirective) || inTexDirective)
-	{
-		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)
-{
-	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)
-{
-
-	decudaInst *instEntry = new decudaInst();
-	instEntry->setBase(".target");
-	g_headerList->add(instEntry);	
-
-	g_headerList->getListEnd().addOperand(firstTarget);
-}
-
-void target_header2(char* firstTarget, char* 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)
-{
-        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 )
-{
-	inEntryDirective = true;
-}
-
-void* reset_symtab()
-{
-	inEntryDirective = false;
-	void* a;
-	return a;
-}
-
-void func_header(const char* 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)
-{
-	const char* compareString = g_headerList->getListEnd().getBase();
-
-	if(inEntryDirective) {
-		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;
-		}
-	}
-
-	if(inTexDirective) {
-		inTexDirective = false;
-		g_headerList->addTex(headerInfo);
-	}
-
-}
diff --git a/setup_environment b/setup_environment
index 6420110..4bda865 100644
--- a/setup_environment
+++ b/setup_environment
@@ -17,9 +17,6 @@ fi
 export GPGPUSIM_ROOT="$( cd "$( dirname "$BASH_SOURCE" )" && pwd )"
 
 export QTINC=/usr/include
-export BOOST_ROOT=/usr/include/boost
-export BOOST_LIB=/usr/lib64
-export BOOST_VER=
 
 # change NVOPENCL_LIBDIR to point to your opencl library directory, usually
 # /usr/lib or /usr/lib64. Not setting this variable will cause gpgpu-sim to