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]

15 files changed, 3 insertions, 4476 deletions

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/header.l b/decuda_to_ptxplus/header.l
deleted file mode 100644
index 5def3bd..0000000
--- a/decuda_to_ptxplus/header.l
+++ /dev/null
@@ -1,124 +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.
-*/
-
-%option noyywrap
-%option yylineno
-%option prefix="header_"
-%{
-#include <iostream>
-#include "headerParser.hpp"
-#include <string.h>
-
-/* not sure if this is needed*/
-#define	YY_NEVER_INTERACTIVE	1
-
-void header_error(const char*);
-extern void output(const char * text);
-%}
-
-
-/*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})
-anycharbutrightbrace [^"}"]
-
-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?
-
-%%
-	/*Translation rules*/
-
-	/*Compiler directives*/
-\.version	return DOTVERSION;
-\.target	return DOTTARGET;
-\.entry	return DOTENTRY;
-
-	/*State spaces*/
-\.param	return DOTPARAM;
-
-	/*operand types*/
-\.u64	return DOTU64;
-\.u32	return DOTU32;
-\.u16	return DOTU16;
-\.b32	return DOTB32;
-\.f32	return DOTF32;
-
-
-	/*identifier + literals*/
-{identifier}	header_lval.string_value = strdup(yytext); return IDENTIFER;
-{decliteral}	header_lval.string_value = strdup(yytext); return DECLITERAL; //change to int later?
-
-	/*braces and parenthesis*/
-\(	return LEFTPAREN;
-\)	return RIGHTPAREN;
-
-	/*discard newline*/
-{newline}+
-
-
-	/*discard commas and periods*/
-\,
-\.
-	/*discard whitespace*/
-{whitespace}+
-	/*discard comments*/
-"//"{anycharbutnewline}*{newline}
-
-	/*discard other header stuff*/
-\.file{anycharbutnewline}*{newline}
-\.extern{anycharbutnewline}*{newline}
-
-	/*discard kernel body*/
-\{{anycharbutrightbrace}*\}
-
-	/*reached end of file*/
-<<EOF>>	return 0;
-
-	/*None of the rules matched. Throw error*/
-.			header_error("Invalid token");
-
-%%
-
-void header_error(const char* message)
-{
-	output("   "); output(message); output(" near \""); output(header_text); output("\"");
-	output(" on line "); 
-	char line[5]; sprintf(line, "%i", header_lineno); output(line);
-	output("\n");
-}
diff --git a/decuda_to_ptxplus/header.y b/decuda_to_ptxplus/header.y
deleted file mode 100644
index 6aa545e..0000000
--- a/decuda_to_ptxplus/header.y
+++ /dev/null
@@ -1,116 +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_headerList;
-extern void output(const char * text);
-%}
-
-
-%union {
-  double double_value;
-  float  float_value;
-  int    int_value;
-  char * string_value;
-  void * ptr_value;
-}
-
-%token DOTVERSION
-%token DOTTARGET
-%token DOTENTRY
-
-%token DOTPARAM
-
-%token DOTU64
-%token DOTU32
-%token DOTU16
-%token DOTB32
-%token DOTF32
-
-%token <string_value> IDENTIFER
-
-	/*change these 4 to int later?*/
-%token <string_value> DECLITERAL
-
-%token LEFTPAREN
-%token RIGHTPAREN
-
-
-%%
-
-	/*translation rules*/
-program		: statementList			{ output("No parsing errors\n");  }
-		;
-
-statementList	: statementList statement	{ output("\n"); }
-		| statement			{ output("\n"); }
-		;
-
-statement	: compilerDirective literal literal	{}
-		| compilerDirective identifierList		{}
-		| 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);}
-			;
-
-identifierList	: identifierList IDENTIFER	{ output(" "); output($2); g_headerList->getListEnd().addOperand($2); }
-		| IDENTIFER			{ output(" "); output($1); g_headerList->getListEnd().addOperand($1); }
-		;
-
-parameterList	: parameterList parameter
-		| 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); }
-		;
-
-opTypes		: DOTU64		{ output(".u64"); g_headerList->getListEnd().addTypeModifier(".u64");}
-		| DOTU32		{ output(".u32"); g_headerList->getListEnd().addTypeModifier(".u32");}
-		| DOTU16		{ output(".u16"); g_headerList->getListEnd().addTypeModifier(".u16");}
-		| DOTB32		{ output(".b32"); g_headerList->getListEnd().addTypeModifier(".b32");}
-		| DOTF32		{ output(".f32"); g_headerList->getListEnd().addTypeModifier(".f32");}
-		;
-
-literal		: DECLITERAL	{ output(" "); output($1); g_headerList->getListEnd().addOperand($1); }
-		;
-
-%%
-
-/*support c++ functions go here*/
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/decuda_to_ptxplus/stringList.cc b/decuda_to_ptxplus/stringList.cc
deleted file mode 100644
index 099e2b9..0000000
--- a/decuda_to_ptxplus/stringList.cc
+++ /dev/null
@@ -1,128 +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 "stringList.h"
-#include <stdio.h>
-#include <iostream>
-
-//Constructor
-stringList::stringList()
-{
-	//initilize everything to empty
-	m_size = 0;
-	m_listStart = NULL;
-	m_listEnd = NULL;
-}
-
-//retrieve list size
-int stringList::getSize()
-{
-	return m_size;
-}
-
-//retrieve pointer to list start
-stringListPiece* stringList::getListStart()
-{
-	return m_listStart;
-}
-
-//retrieve point to list end
-stringListPiece* stringList::getListEnd()
-{
-	return m_listEnd;
-}
-
-//add string to end of the list
-//initilize list starting point if this is the first entry
-//increment list size count
-int stringList::add(stringListPiece* newString)
-{
-
-	if(m_listStart==NULL)
-		m_listStart=newString;
-	else
-		m_listEnd->nextString = newString;
-
-	m_listEnd = newString;
-	m_listEnd->nextString = NULL;
-
-	return m_size++;
-}
-
-bool stringList::remove(int index)
-{
-	if(index >= m_size ) return false;
-
-	stringListPiece* m_remove;
-	stringListPiece* currentPiece;
-
-	if(m_size == 1) {
-		m_remove = m_listStart;
-		m_listStart = NULL;
-		m_listEnd = NULL;
-	} else {
-
-		if(index == 0) {
-			m_remove = m_listStart;
-			m_listStart = m_remove->nextString;
-		} else if(index == m_size - 1) {
-			m_remove = m_listEnd;
-			currentPiece = m_listStart;
-			for(int i=1; i<m_size-1; i++)
-			{
-				currentPiece = currentPiece->nextString;
-			}
-			currentPiece->nextString = NULL;
-			m_listEnd = currentPiece;
-		} else {
-			currentPiece = m_listStart;
-			for(int i=1; i<=index-1; i++)
-			{
-				currentPiece = currentPiece->nextString;
-			}
-			m_remove = currentPiece->nextString;
-			currentPiece->nextString = m_remove->nextString;
-		}
-	}
-
-	delete m_remove;
-	m_size -= 1;
-	return true;
-}
-
-//print out all the Decuda Instructions in the list
-void stringList::printStringList()
-{
-	stringListPiece* currentPiece = m_listStart;
-
-	for(int i=0; (i<m_size)&&(currentPiece!=NULL); i++)
-	{
-		std::cout << currentPiece->stringText << " ";
-		currentPiece = currentPiece->nextString;
-	}
-}
diff --git a/decuda_to_ptxplus/stringList.h b/decuda_to_ptxplus/stringList.h
deleted file mode 100644
index e1dfaf0..0000000
--- a/decuda_to_ptxplus/stringList.h
+++ /dev/null
@@ -1,62 +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.
-
-
-struct stringListPiece
-{
-	const char* stringText;
-	
-	stringListPiece* nextString;
-};
-
-
-class stringList
-{
-
-private:
-	//List size
-	int m_size;
-
-	//Start is the first entry, end is the last entry
-	stringListPiece* m_listStart;
-	stringListPiece* m_listEnd;
-public:
-	//constructor
-	stringList();
-
-	//accessors
-	int getSize();
-	stringListPiece* getListStart();
-	stringListPiece* getListEnd();
-
-	//mutator
-	int add(stringListPiece* newString); //add String to list
-	bool remove(int index); //remove string at index
-
-	//print representation
-	void printStringList();
-};