summaryrefslogtreecommitdiff
path: root/decuda_to_ptxplus
diff options
context:
space:
mode:
authorJimmy Kwa <[email protected]>2010-10-28 12:07:29 -0800
committerJimmy Kwa <[email protected]>2010-10-28 12:07:29 -0800
commit0586c55d7688f492d5264cb9299cf00ff73d3df7 (patch)
treeef1ecc58933398e333f16dc56915bdffc64ef422 /decuda_to_ptxplus
parent0efd3c00f5611bfa82b01d87d175122388d621cc (diff)
used p4 integrate to copy decuda_to_ptxplus into the distribtion folder. Edited baseline_ptxplus.config so the ptxplus regression runs properly.
[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7942]
Diffstat (limited to 'decuda_to_ptxplus')
-rw-r--r--decuda_to_ptxplus/decuda.l273
-rw-r--r--decuda_to_ptxplus/decuda.y307
-rw-r--r--decuda_to_ptxplus/decudaInst.cc1473
-rw-r--r--decuda_to_ptxplus/decudaInst.h66
-rw-r--r--decuda_to_ptxplus/decudaInstList.cc761
-rw-r--r--decuda_to_ptxplus/decudaInstList.h146
-rw-r--r--decuda_to_ptxplus/decuda_to_ptxplus.cc87
-rw-r--r--decuda_to_ptxplus/header.l95
-rw-r--r--decuda_to_ptxplus/header.y88
-rw-r--r--decuda_to_ptxplus/makefile68
-rw-r--r--decuda_to_ptxplus/opcodes.def159
-rw-r--r--decuda_to_ptxplus/opcodes.h97
-rw-r--r--decuda_to_ptxplus/ptx_parser.h283
-rw-r--r--decuda_to_ptxplus/stringList.cc100
-rw-r--r--decuda_to_ptxplus/stringList.h34
15 files changed, 4037 insertions, 0 deletions
diff --git a/decuda_to_ptxplus/decuda.l b/decuda_to_ptxplus/decuda.l
new file mode 100644
index 0000000..e47a4c7
--- /dev/null
+++ b/decuda_to_ptxplus/decuda.l
@@ -0,0 +1,273 @@
+%{
+#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;
+\.lo yylval.string_value = strdup(yytext); return DOTLO;
+\.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
new file mode 100644
index 0000000..52065d2
--- /dev/null
+++ b/decuda_to_ptxplus/decuda.y
@@ -0,0 +1,307 @@
+%{
+#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 { $$=".u32"}
+ | DDOTF32 { $$=".f32"}
+ ;
+
+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
new file mode 100644
index 0000000..3e5a402
--- /dev/null
+++ b/decuda_to_ptxplus/decudaInst.cc
@@ -0,0 +1,1473 @@
+#include "decudaInst.h"
+#include <iostream>
+#include <string>
+#include <cstring>
+#include <stdlib.h>
+
+extern void output(const char * text);
+
+//Constructor
+decudaInst::decudaInst()
+{
+ //initilize everything to empty
+ m_label = "";
+ m_predicate = "";
+ 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
+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(char* setBaseValue)
+{
+ m_base = setBaseValue;
+}
+
+void decudaInst::addBaseModifier(char* addBaseMod)
+{
+ stringListPiece* tempPiece = new stringListPiece;
+ tempPiece->stringText = addBaseMod;
+
+ m_baseModifiers->add(tempPiece);
+}
+
+void decudaInst::addTypeModifier(char* addTypeMod)
+{
+ stringListPiece* tempPiece = new stringListPiece;
+ tempPiece->stringText = addTypeMod;
+
+ m_typeModifiers->add(tempPiece);
+}
+
+void decudaInst::addOperand(char* addOp)
+{
+ stringListPiece* tempPiece = new stringListPiece;
+ tempPiece->stringText = addOp;
+
+ m_operands->add(tempPiece);
+}
+
+void decudaInst::setPredicate(char* setPredicateValue)
+{
+ m_predicate = setPredicateValue;
+}
+
+void decudaInst::addPredicateModifier(char* addPredicateMod)
+{
+ stringListPiece* tempPiece = new stringListPiece;
+ tempPiece->stringText = addPredicateMod;
+
+ m_predicateModifiers->add(tempPiece);
+}
+
+void decudaInst::setLabel(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()
+{
+ if(m_predicate != "") {
+ output(m_predicate);
+
+ stringListPiece* currentPiece = m_predicateModifiers->getListStart();
+ for(int i=0; (i<m_predicateModifiers->getSize())&&(currentPiece!=NULL); i++)
+ {
+ output(currentPiece->stringText);
+
+ currentPiece = currentPiece->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();
+ 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 = 0; //0=default, 1=bb64 type, 2=bb128 type
+
+ // Replace '%clock' with '%halfclock'
+ currentPiece = m_operands->getListStart();
+ for(int i=0; (i<m_operands->getSize())&&(currentPiece!=NULL); i++)
+ {
+ char* modString = currentPiece->stringText;
+ if( strcmp(modString, "%%clock")==0 ) {
+ 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++)
+ {
+ 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++)
+ {
+ char* modString = currentPiece->stringText;
+ if( strcmp(modString, ".end")==0 ) {
+ 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
+ currentPiece = m_typeModifiers->getListStart();
+ for(int i=0; (i<m_typeModifiers->getSize())&&(currentPiece!=NULL); i++)
+ {
+ char* modString = currentPiece->stringText;
+ if( strcmp(modString, ".b64")==0 ) {
+ char* newText = ".bb64";
+ currentPiece->stringText = newText;
+ vectorFlag = 1;
+ }
+ if( strcmp(modString, ".b128")==0 ) {
+ char* newText = ".bb128";
+ currentPiece->stringText = newText;
+ vectorFlag = 2;
+ }
+ currentPiece = currentPiece->nextString;
+ }
+
+ //TODO: add support for two vectors in one instruction. Currently not sure if it can actually occur or not
+ if(vectorFlag != 0)
+ {
+ currentPiece = m_operands->getListStart();
+ for(int i=0; (i<m_operands->getSize())&&(currentPiece!=NULL); i++)
+ {
+ char newText[40] = "";
+ char *regNumString;
+ int regNumInt;
+
+ char* modString = currentPiece->stringText;
+ if( modString[0] == '$' && modString[1] == 'r' ) {
+ strcpy(newText, modString);
+ strtok (newText, "r");
+ regNumString = strtok (NULL, "r");
+ regNumInt = atoi(regNumString);
+ if(vectorFlag ==1)
+ snprintf(newText,40,"{$r%u,$r%u}", regNumInt+0, regNumInt+1);
+ if(vectorFlag ==2)
+ snprintf(newText,40,"{$r%u,$r%u,$r%u,$r%u}", regNumInt+0, regNumInt+1, regNumInt+2, regNumInt+3);
+
+ currentPiece->stringText = newText;
+ break;
+ } else if( modString[0] == '$' && modString[1] == 'o'&& modString[2] == 'f' ) {
+ strcpy(newText, modString);
+ strtok (newText, "s");
+ regNumString = strtok (NULL, "s");
+ regNumInt = atoi(regNumString);
+ if(vectorFlag ==1)
+ snprintf(newText,40,"{$ofs%u,$ofs%u}", regNumInt+0, regNumInt+1);
+ if(vectorFlag ==2)
+ snprintf(newText,40,"{$ofs%u,$ofs%u,$ofs%u,$ofs%u}", regNumInt+0, regNumInt+1, regNumInt+2, regNumInt+3);
+
+ currentPiece->stringText = newText;
+ break;
+ } else if( modString[0] == '$' && modString[1] == 'p' ) {
+ strcpy(newText, modString);
+ strtok (newText, "p");
+ regNumString = strtok (NULL, "p");
+ regNumInt = atoi(regNumString);
+ if(vectorFlag ==1)
+ snprintf(newText,40,"{$p%u,$p%u}", regNumInt+0, regNumInt+1);
+ if(vectorFlag ==2)
+ snprintf(newText,40,"{$p%u,$p%u,$p%u,$p%u}", regNumInt+0, regNumInt+1, regNumInt+2, regNumInt+3);
+
+ currentPiece->stringText = newText;
+ break;
+ } else if( modString[0] == '$' && modString[1] == 'o' ) {
+ strcpy(newText, modString);
+ strtok (newText, "o");
+ regNumString = strtok (NULL, "o");
+ regNumInt = atoi(regNumString);
+ if(vectorFlag ==1)
+ snprintf(newText,40,"{$o%u,$o%u}", regNumInt+0, regNumInt+1);
+ if(vectorFlag ==2)
+ snprintf(newText,40,"{$o%u,$o%u,$o%u,$o%u}", regNumInt+0, regNumInt+1, regNumInt+2, regNumInt+3);
+
+ currentPiece->stringText = newText;
+ break;
+ } else if( modString[0] == '$' && modString[1] == 't'&& modString[2] == 'e' ) {
+ strcpy(newText, modString);
+ strtok (newText, "x");
+ regNumString = strtok (NULL, "x");
+ regNumInt = atoi(regNumString);
+ if(vectorFlag ==1)
+ snprintf(newText,40,"{$tex%u,$tex%u}", regNumInt+0, regNumInt+1);
+ if(vectorFlag ==2)
+ snprintf(newText,40,"{$tex%u,$tex%u,$tex%u,$tex%u}", regNumInt+0, regNumInt+1, regNumInt+2, regNumInt+3);
+
+ currentPiece->stringText = newText;
+ break;
+ }
+ currentPiece = currentPiece->nextString;
+ }
+ }
+
+ // Predicate modifiers :
+ // @$p0.ne and @$p0.neu map to @$p0,
+ // $p0.eq, $p0.equ, $p0.ge map to @!$p0
+ // Assuming only one modifier for now, only removes the first one
+ /*if(m_predicate != "")
+ {
+ currentPiece = m_predicateModifiers->getListStart();
+ char* modString = currentPiece->stringText;
+ if( strcmp(modString, ".ne")==0 ||
+ strcmp(modString, ".neu")==0 ||
+ strcmp(modString, ".gtu")==0 ||
+ strcmp(modString, ".ltu")==0 ||
+ strcmp(modString, ".sf")==0 ) {
+ m_predicateModifiers->remove(0);
+ } else if( strcmp(modString, ".eq")==0 ||
+ strcmp(modString, ".equ")==0 ||
+ strcmp(modString, ".le")==0 ||
+ strcmp(modString, ".ge")==0 ||
+ strcmp(modString, ".nsf")==0 ) {
+ // Insert '!' into predicate
+ std::string temp = m_predicate;
+ temp.insert(1, "!");
+ strcpy(m_predicate, temp.c_str());
+ m_predicateModifiers->remove(0);
+ }
+ }*/
+
+ //
+ // 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++)
+ {
+ 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");
+
+ char* type = m_typeModifiers->getListStart()->stringText;
+ output(type);
+
+ printOperands();
+
+ output(";");
+
+ } else {
+ // cvt instruction
+
+ // Remove .neg base modifier
+ /*currentPiece = m_baseModifiers->getListStart();
+ for(int i=0; (i<m_baseModifiers->getSize())&&(currentPiece!=NULL); i++)
+ {
+ char* modString = currentPiece->stringText;
+ if( strcmp(modString, ".neg")==0 ) {
+ m_baseModifiers->remove(i);
+ }
+ currentPiece = currentPiece->nextString;
+ }*/
+
+ printLabel();
+ printPredicate();
+
+ output("cvt");
+
+ int typeModifiers = m_typeModifiers->getSize();
+
+ if(typeModifiers == 2) {
+ if(cvt_neg_mod_flag == 0)
+ stringListPiece* currentPiece = m_baseModifiers->getListStart();
+ else
+ stringListPiece* currentPiece = m_baseModifiers->getListStart()->nextString;
+
+ char* dstType = m_typeModifiers->getListStart()->stringText;
+ char* srcType = m_typeModifiers->getListStart()->nextString->stringText;
+
+ for(int i=0; (i<m_baseModifiers->getSize())&&(currentPiece!=NULL); i++)
+ {
+ 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, ".f64")==0))
+ output(modString);
+ if((strcmp(dstType, ".f16")==0) && (strcmp(srcType, ".f64")==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, ".f64")==0) && (strcmp(srcType, ".u8")==0))
+ output(modString);
+ if((strcmp(dstType, ".f64")==0) && (strcmp(srcType, ".u16")==0))
+ output(modString);
+ if((strcmp(dstType, ".f64")==0) && (strcmp(srcType, ".u32")==0))
+ output(modString);
+ if((strcmp(dstType, ".f64")==0) && (strcmp(srcType, ".u64")==0))
+ output(modString);
+ if((strcmp(dstType, ".f64")==0) && (strcmp(srcType, ".s8")==0))
+ output(modString);
+ if((strcmp(dstType, ".f64")==0) && (strcmp(srcType, ".s16")==0))
+ output(modString);
+ if((strcmp(dstType, ".f64")==0) && (strcmp(srcType, ".s32")==0))
+ output(modString);
+ if((strcmp(dstType, ".f64")==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, ".f64")==0))
+ { output(modString); output("i"); }
+ if((strcmp(dstType, ".u16")==0) && (strcmp(srcType, ".f64")==0))
+ { output(modString); output("i"); }
+ if((strcmp(dstType, ".u32")==0) && (strcmp(srcType, ".f64")==0))
+ { output(modString); output("i"); }
+ if((strcmp(dstType, ".u64")==0) && (strcmp(srcType, ".f64")==0))
+ { output(modString); output("i"); }
+ if((strcmp(dstType, ".s8")==0) && (strcmp(srcType, ".f64")==0))
+ { output(modString); output("i"); }
+ if((strcmp(dstType, ".s16")==0) && (strcmp(srcType, ".f64")==0))
+ { output(modString); output("i"); }
+ if((strcmp(dstType, ".s32")==0) && (strcmp(srcType, ".f64")==0))
+ { output(modString); output("i"); }
+ if((strcmp(dstType, ".s64")==0) && (strcmp(srcType, ".f64")==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, ".f64")==0) && (strcmp(srcType, ".f64")==0))
+ { output(modString); output("i"); }
+ }
+ else if (
+ (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
+ 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
+ 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
+ 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);
+ }
+ char* firstOperand = m_operands->getListStart()->stringText;
+ char* secondOperand = m_operands->getListStart()->nextString->stringText;
+ 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
new file mode 100644
index 0000000..a1f0a8e
--- /dev/null
+++ b/decuda_to_ptxplus/decudaInst.h
@@ -0,0 +1,66 @@
+#include "stringList.h"
+#include <assert.h>
+
+class decudaInst
+{
+
+private:
+ //instruction data
+ char* m_label; //instruction label
+ char* m_predicate; //instruction predicate
+ char* m_base; //instruction mnemonic
+ stringList* m_baseModifiers; //base modifiers
+ stringList* m_typeModifiers; //operand types
+ stringList* m_operands; //operands
+ stringList* m_predicateModifiers; //predicate modifiers
+
+
+ int m_opPerCycle; // stores operations per cycle count for current instruction
+
+ //next instruction in linked list
+ //direction is m_listStart to m_listEnd
+ decudaInst* m_nextDecudaInst;
+
+ // print instruction unmodified
+ void printDefaultPtx();
+ void printBaseModifiers();
+ void printTypeModifiers();
+ void printOperands();
+ void printLabel();
+ void printPredicate();
+
+public:
+ //constructor
+ decudaInst();
+
+ //accessors
+ char* getBase();
+ stringList* getOperands();
+ stringList* getBaseModifiers();
+ stringList* getTypeModifiers();
+ decudaInst* getNextDecudaInst();
+
+ int getOpPerCycle() const { return m_opPerCycle; }
+
+ bool isEntryStart(); // true if start of an entry
+
+ //mutators
+ void setBase(char* setBaseValue);
+ void addBaseModifier(char* addBaseMod);
+ void addTypeModifier(char* addTypeMod);
+ void addOperand(char* addOp);
+ void setPredicate(char* setPredicateValue);
+ void addPredicateModifier(char* addPredicateMod);
+ void setLabel(char* setLabelValue);
+
+ void setNextDecudaInst(decudaInst* setDecudaInstValue);
+
+ //print representation
+ bool printHeaderInst();
+ bool printHeaderInst2();
+ void printDecudaInst();
+ void printNewPtx();
+ void printHeaderPtx();
+ //TODO: translate to New PTX and print
+
+};
diff --git a/decuda_to_ptxplus/decudaInstList.cc b/decuda_to_ptxplus/decudaInstList.cc
new file mode 100644
index 0000000..4605e64
--- /dev/null
+++ b/decuda_to_ptxplus/decudaInstList.cc
@@ -0,0 +1,761 @@
+#include "decudaInstList.h"
+
+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;
+ }
+ }
+}
+
+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++)
+ {
+ char* modString = currentPiece->stringText;
+
+ if( strcmp(modString, ".b64")==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())];
+ 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())];
+ 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())];
+ strcpy(doublePredRegName, doublePredReg.c_str());
+ getListEnd().addOperand(doublePredRegName);
+}
+
+
+// add to tex list
+void decudaInstList::addTex(std::string tex)
+{
+ std::string origTex = tex;
+
+ // 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())];
+ 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);
+}
+
+// set type of constant memory
+void decudaInstList::setConstMemoryType(char* type)
+{
+ m_constMemoryList.back().type = type;
+}
+
+// print const memory directive
+void decudaInstList::printMemory()
+{
+
+ // Constant memory
+ std::list<constMemory>::iterator i;
+ for(i=m_constMemoryList.begin(); i!=m_constMemoryList.end(); ++i) {
+ char line[40];
+
+ // Global or entry specific
+ if(i->entryIndex == 0)
+ sprintf(line, ".const %s c%d[%d] = {", i->type, i->index, 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");
+ }
+
+ // 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())];
+ 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;
+
+ // 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 = 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())];
+ 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
new file mode 100644
index 0000000..649a5f5
--- /dev/null
+++ b/decuda_to_ptxplus/decudaInstList.h
@@ -0,0 +1,146 @@
+#include "decudaInst.h"
+#include <list>
+#include <string>
+#include <cstring>
+#include <stdio.h>
+#include <iostream>
+#include <sstream>
+#include <stdlib.h>
+#include <regex.hpp>
+#include <map>
+
+// Used for entry specific constant memory segments (c1)
+struct constMemory
+{
+ int index;
+ int entryIndex;
+ 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;
+
+ // Histogram for operation per cycle count
+ std::map<std::string, int> m_opPerCycleHistogram;
+};
+
+
+
+class decudaInstList
+{
+
+private:
+ // List of decuda entries
+ std::list<decudaEntry> m_entryList;
+
+ // Const memory list
+ std::list<constMemory> m_constMemoryList;
+
+ // 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(char* type); // set type of constant memory
+ void addConstMemoryValue(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
new file mode 100644
index 0000000..b7d7132
--- /dev/null
+++ b/decuda_to_ptxplus/decuda_to_ptxplus.cc
@@ -0,0 +1,87 @@
+#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
new file mode 100644
index 0000000..848c189
--- /dev/null
+++ b/decuda_to_ptxplus/header.l
@@ -0,0 +1,95 @@
+%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
new file mode 100644
index 0000000..c0554f3
--- /dev/null
+++ b/decuda_to_ptxplus/header.y
@@ -0,0 +1,88 @@
+%{
+#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
new file mode 100644
index 0000000..cc0a37e
--- /dev/null
+++ b/decuda_to_ptxplus/makefile
@@ -0,0 +1,68 @@
+FLAGS = -ggdb -I${BOOST_ROOT} -lboost_regex -lboost_filesystem
+
+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 : ../distribution/src/cuda-sim/ptx.l
+ flex ../distribution/src/cuda-sim/ptx.l
+
+ptx.tab.c : ../distribution/src/cuda-sim/ptx.y
+ bison --name-prefix=ptx_ -v -d ../distribution/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
+
+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
diff --git a/decuda_to_ptxplus/opcodes.def b/decuda_to_ptxplus/opcodes.def
new file mode 100644
index 0000000..7815f24
--- /dev/null
+++ b/decuda_to_ptxplus/opcodes.def
@@ -0,0 +1,159 @@
+/*
+ * 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
new file mode 100644
index 0000000..2ceafaf
--- /dev/null
+++ b/decuda_to_ptxplus/opcodes.h
@@ -0,0 +1,97 @@
+/*
+ * 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
new file mode 100644
index 0000000..c0b31b8
--- /dev/null
+++ b/decuda_to_ptxplus/ptx_parser.h
@@ -0,0 +1,283 @@
+ #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
+
+
+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(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;
+ 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 )
+{
+ switch(spec)
+ {
+ case param_space_unclassified:
+ if(inEntryDirective && inParamDirective) {
+ decudaInst *instEntry = new decudaInst();
+ instEntry->setBase(".param");
+ g_headerList->add(instEntry);
+ }
+ break;
+ case tex_space:
+ inTexDirective = true;
+ decudaInst *instEntry = new decudaInst();
+ instEntry->setBase(".tex");
+ g_headerList->add(instEntry);
+ break;
+ }
+}
+
+void add_scalar_type_spec( int headerInput )
+{
+ 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(char* headerBase)
+{
+
+ // If start of an entry
+ if((strcmp(headerBase, ".entry")==0)||(strcmp(headerBase, ".func")==0)) {
+ g_headerList->addEntry("");
+
+ decudaInst *instEntry = new decudaInst();
+ instEntry->setBase(headerBase);
+ g_headerList->add(instEntry);
+
+ }
+}
+
+void func_header_info(char* headerInfo)
+{
+ 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
new file mode 100644
index 0000000..680a860
--- /dev/null
+++ b/decuda_to_ptxplus/stringList.cc
@@ -0,0 +1,100 @@
+#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
new file mode 100644
index 0000000..25d10f7
--- /dev/null
+++ b/decuda_to_ptxplus/stringList.h
@@ -0,0 +1,34 @@
+struct stringListPiece
+{
+ 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();
+};
generated by cgit v1.3 (git 2.53.0) at 2026-07-10 23:40:40 +0000