summaryrefslogtreecommitdiff
path: root/decuda_to_ptxplus/decudaInst.cc
diff options
context:
space:
mode:
Diffstat (limited to 'decuda_to_ptxplus/decudaInst.cc')
-rw-r--r--decuda_to_ptxplus/decudaInst.cc1473
1 files changed, 1473 insertions, 0 deletions
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);
+ }
+
+}
generated by cgit v1.3 (git 2.53.0) at 2026-07-11 00:52:33 +0000