diff options
Diffstat (limited to 'cuobjdump_to_ptxplus/cuobjdumpInst.cc')
| -rw-r--r-- | cuobjdump_to_ptxplus/cuobjdumpInst.cc | 1690 |
1 files changed, 1690 insertions, 0 deletions
diff --git a/cuobjdump_to_ptxplus/cuobjdumpInst.cc b/cuobjdump_to_ptxplus/cuobjdumpInst.cc new file mode 100644 index 0000000..b014888 --- /dev/null +++ b/cuobjdump_to_ptxplus/cuobjdumpInst.cc @@ -0,0 +1,1690 @@ +#include "cuobjdumpInst.h" +#include <sstream> +#include <iostream> + +extern void output(const char * text); + +//Constructor +cuobjdumpInst::cuobjdumpInst() +{ + //initilize everything to empty +} + +bool cuobjdumpInst::checkCubojdumpLabel(std::list<std::string> labelList, std::string label) +{ + if(labelList.empty()) + return false; + + std::list<std::string>::iterator labelIterator; + + for( labelIterator=labelList.begin(); labelIterator!=labelList.end(); labelIterator++ ) + { + if(label.compare(*labelIterator) == 0) + return true; + } + + return false; +} + +void cuobjdumpInst::printCuobjdumpLabel(std::list<std::string> labelList) +{ + if((strcmp(m_label, "")!=0)&&(checkCubojdumpLabel(labelList, m_label))) { + output(m_label); + output(": "); + } +} + +void cuobjdumpInst::printCuobjdumpPredicate() +{ + stringListPiece* currentPiece = m_predicate->getListStart(); + if(currentPiece!=NULL) + { + output("@$p"); + char modString[2]; + modString[0]=currentPiece->stringText[1]; + modString[1]='\0'; + output(modString); + + stringListPiece* currentPiece2 = m_predicateModifiers->getListStart(); + for(int i=0; (i<m_predicateModifiers->getSize())&&(currentPiece2!=NULL); i++) + { + const char* modString2 = currentPiece2->stringText; + char* modString3 = new char[strlen(currentPiece2->stringText)+1]; + + for(int i=0; i<strlen(modString2); i++) + { + if(modString2[i] >= 'A' && modString2[i] <= 'Z') + modString3[i] = modString2[i] + 32; + else + modString3[i] = modString2[i]; + } + modString3[strlen(modString2)]='\0'; + if( strcmp(modString3, ".not_sign")==0) + { + output(".nsf"); + } else if( strcmp(modString3, ".sign")==0) + { + output(".sf"); + } else if( strcmp(modString3, ".carry")==0) + { + output(".cf"); + } + else + { + output(modString3); + } + currentPiece = currentPiece2->nextString; + } + output(" "); + } +} + +void cuobjdumpInst::printCuobjdumpTypeModifiers() +{ + stringListPiece* currentPiece = m_typeModifiers->getListStart(); + for(int i=0; (i<m_typeModifiers->getSize())&&(currentPiece!=NULL); i++) + { + const char* modString = currentPiece->stringText; + if( strcmp(modString, ".F16")==0 ) + output(".f16"); + else if( strcmp(modString, ".F32")==0 ) + output(".f32"); + else if( strcmp(modString, ".F64")==0 ) + output(".ff64"); + else if( strcmp(modString, ".S8")==0 ) + output(".s8"); + else if( strcmp(modString, ".S16")==0 ) + output(".s16"); + else if( strcmp(modString, ".S32")==0 ) + output(".s32"); + else if( strcmp(modString, ".S64")==0 ) + output(".bb64"); //TODO: might have to change to .ss64 in the future. + else if( strcmp(modString, ".S128")==0 ) + output(".bb128"); //TODO: might have to change to .ss64 in the future. + else if( strcmp(modString, ".U8")==0 ) + output(".u8"); + else if( strcmp(modString, ".U16")==0 ) + output(".u16"); + else if( strcmp(modString, ".U32")==0 ) + output(".u32"); + else if( strcmp(modString, ".U64")==0 ) + output(".bb64"); //TODO: might have to change to .ss64 in the future. + else if( strcmp(modString, ".HI")==0 ) + output(".hi"); + else + { + printf("Unknown Type: "); + printf(modString); + printf("\n"); + output("Unknown Type: "); output(modString); + assert(0); + } + + currentPiece = currentPiece->nextString; + } +} + +void cuobjdumpInst::printCuobjdumpBaseModifiers() +{ + stringListPiece* currentPiece = m_baseModifiers->getListStart(); + for(int i=0; (i<m_baseModifiers->getSize())&&(currentPiece!=NULL); i++) + { + const char* modString = currentPiece->stringText; + if( strcmp(modString, "EQ")==0 ) + output(".eq"); + else if( strcmp(modString, "EQU")==0 ) + output(".equ"); + else if( strcmp(modString, "GE")==0 ) + output(".ge"); + else if( strcmp(modString, "GEU")==0 ) + output(".geu"); + else if( strcmp(modString, "GT")==0 ) + output(".gt"); + else if( strcmp(modString, "GTU")==0 ) + output(".gtu"); + else if( strcmp(modString, "LE")==0 ) + output(".le"); + else if( strcmp(modString, "LEU")==0 ) + output(".leu"); + else if( strcmp(modString, "LT")==0 ) + output(".lt"); + else if( strcmp(modString, "LTU")==0 ) + output(".ltu"); + else if( strcmp(modString, "NE")==0 ) + output(".ne"); + else if( strcmp(modString, "NEU")==0 ) + output(".neu"); + else if( strcmp(modString, ".exit")==0 ) + output(modString); + else if( strcmp(modString, ".abs")==0 ) + { + if((strcmp(m_base, "F2F")!=0) && (strcmp(m_base, "I2I")!=0)) + { + output(modString); + } + } + else if( strcmp(modString, "ex2")==0 ) + output(modString); + else if( strcmp(modString, "sin")==0 ) + output(modString); + else if( strcmp(modString, "cos")==0 ) + output(modString); + else if( strcmp(modString, ".rz")==0 ) + output(modString); + else if( strcmp(modString, ".bext")==0 ) + { + //".bext" is a modifier that indicated u8 to u16 type conversion, I think + } + else if( strcmp(modString, ".s")==0 ) + { + //".s" is the same as ".join" in decuda. + } + else if( strcmp(modString, ".sfu")==0 ) + { + //".sfu" is an unknown base modifier, TODO: find out what it is + } + else if( strcmp(modString, ".x")==0 ) + { + //".x" is an unknown base modifier, TODO: find out what it is + output(modString); + } + else if( strcmp(modString, ".e")==0 ) + { + //".e" is an unknown base modifier, TODO: find out what it is + output(modString); + } + else if( strcmp(modString, ".ir")==0 ) + { + //".ir" is an unknown base modifier, TODO: find out what it is + output(modString); + } + else if(strcmp(modString, "IADD")==0 || + strcmp(modString, "IMIN")==0 || + strcmp(modString, "IMAX")==0 + //strcmp(modString, "INC")==0 || + //strcmp(modString, "DEC")==0 || + //strcmp(modString, "IAND")==0 || + //strcmp(modString, "IOR")==0 || + //strcmp(modString, "IXOR")==0 + ) + { + /* + * This is the case of a GRED or GATOM operation + */ + output("."); + std::string modstr = modString; + modstr = modstr.substr(1); + for (int i=0; i<modstr.length(); i++){ + modstr[i] = tolower(modstr[i]); + } + output(modstr.c_str()); + } + else + { + printf("Unknown Base Mod: "); + printf(modString); + printf("\n"); + output("Unknown Base Mod: "); + output(modString); + assert(0); + } + + currentPiece = currentPiece->nextString; + } +} + +/* + * Remove the trailing 'l' or 'h' and output the operand followed by ".lo" or ".hi" respectively + */ +void cuobjdumpInst::printCuobjdumpOperandlohi(std::string op) { + if (op.substr(op.length()-1) == "l") { + output(op.substr(0, op.length()-1).c_str()); + output(".lo"); + } else if (op.substr(op.length()-1) == "h"){ + output(op.substr(0, op.length()-1).c_str()); + output(".hi"); + } else { + output(op.c_str()); + } +} + +void cuobjdumpInst::printCuobjdumpOperand(stringListPiece* currentPiece, std::string operandDelimiter, const char* base) +{ + + output(operandDelimiter.c_str()); + output(" "); + + //Current piece + std::string currp = currentPiece->stringText; + std::string mod; + + if(currp[0] == '-') { + mod = currp.substr(1); + + if(mod.substr(0,2) == "0x") { + unsigned int immValue; + std::stringstream hexStringConvert2; + hexStringConvert2 << std::hex << mod; + hexStringConvert2 >> immValue; + + if(immValue){ + immValue = ~immValue + 1; + } + + std::stringstream outputhexstream; + outputhexstream << std::hex << immValue; + std::string outputhex; + outputhex = outputhexstream.str(); + output("0x"); + for(int i=8; i > outputhex.length(); i--) + { + output("0"); + } + output(outputhex.c_str()); + return; + } + + output("-"); + } + else + { + mod = std::string(currp); + } + + // Make it lower case + if(mod.substr(0,9)!= "constant1") + for(int i=0; i<mod.length(); i++) + { + mod[i] = tolower(mod[i]); + } + + //double destination + if(mod[0]=='.') //double destination + { + std::string temp = mod.substr(2,2); + output("$p"); + output(mod.substr(2,2).c_str()); + + mod = mod.substr(4); + } + + + if(mod == "g [0x1].u16") { //handling special register case: %ntid.x + output("%%ntid.x"); + } else if(mod == "g [0x2].u16") { //handling special register case: %ntid.y + output("%%ntid.y"); + } else if(mod == "g [0x3].u16") { //handling special register case: %ntid.x + output("%%ntid.z"); + } else if(mod =="g [0x4].u16") { //handling special register case: %nctaid.x + output("%%nctaid.x"); + } else if(mod =="g [0x5].u16") { //handling special register case: %nctaid.y + output("%%nctaid.y"); + } else if(mod == "g [0x6].u16") {//handling special register case: %ctaid.x + output("%%ctaid.x"); + } else if(mod == "g [0x7].u16") {//handling special register case: %ctaid.y + output("%%ctaid.y"); + } else if(mod[0]=='r') { //basic register + if( (strcmp(m_base, "DADD")==0) || + (strcmp(m_base, "DMUL")==0) || + (strcmp(m_base, "DFMA")==0) || + ( (m_typeModifiers->getSize()==1) && + (strcmp((m_typeModifiers->getListStart()->stringText), ".S64")==0) && + ( (strcmp(m_base, "G2R")==0) || + (strcmp(m_base, "R2G")==0) || + (strcmp(m_base, "GLD")==0) || + (strcmp(m_base, "GST")==0) || + (strcmp(m_base, "LST")==0) ))) { + std::string modsub = mod.substr(1); + int regNumInt = atoi(modsub.c_str()); + std::stringstream temp; + temp << "{$r" << (regNumInt) << ",$r"<< (regNumInt+1) << "}"; + output(temp.str().c_str()); + } else if( (m_typeModifiers->getSize()==1) && + (strcmp((m_typeModifiers->getListStart()->stringText), ".S128")==0)) { + std::string modsub = mod.substr(1); + int regNumInt = atoi(modsub.c_str()); + std::stringstream temp; + temp << "{$r" << (regNumInt); + temp << ",$r" << (regNumInt+1); + temp << ",$r" << (regNumInt+2); + temp << ",$r" << (regNumInt+3) << "}"; + output(temp.str().c_str()); + } else { + output("$"); + printCuobjdumpOperandlohi(mod); + } + } else if(mod[0] == 'c' && mod.length() == 2) { //predicate register (conditional code) + output("$p"); + output(mod.substr(1,1).c_str()); + } else if(mod[0]=='a') {//offset register + output("$ofs"); + mod = mod.substr(1); + printCuobjdumpOperandlohi(mod); + } else if(mod[0]=='o') {//output register + output("$o127"); + } else if ( mod[0]=='g' || + mod.substr(0,2) == "lo" || + mod.substr(0,8) == "constant") { + //memory operands, global14 = global, g = shared, c = constant + std::string modsub; + std::string modsub2; + std::string modsub3; + modsub = mod.c_str(); + int localFlag = 0; + + if(mod.find("global14") != std::string::npos) { + //Those instructions don't need the dereferencing done by g [*] + if( strcmp(base,"GRED")==0 || + strcmp(base, "GATOM")==0 || + strcmp(base, "GST")==0 || + strcmp(base, "GLD")==0) + output("["); + else + output("g["); + } else if(mod[0]=='g') { + //Shared memory + output("s["); + } else if(mod.find("local") != std::string::npos) { + modsub3 = modsub.substr(4, modsub.length()-10); + output(modsub3.c_str()); + output("["); + localFlag = 1; + } else if(mod.substr(0,9) == "constant1") { + output(modsub.substr(0, modsub.find_first_of("[]")+1).c_str()); + } else if(mod.substr(0,9)=="constant0"){ + output("constant0["); + } else { + printf("Unidentified modifier: %s\n", mod.c_str()); + assert(0); + } + + modsub = modsub.substr(modsub.find_first_of("[]")+1); + modsub = modsub.substr(0, modsub.length()-1); + //Here we handle whatever was inside the [] + int plusequalFlag = 0; + if(modsub.find("+++") != std::string::npos) { + plusequalFlag = 1; + } + + if(modsub.find("+") != std::string::npos) + { + //Handle a1+++0x1 or a1+0x1 + modsub2 = modsub.substr(modsub.find("+")+1); + + output("$ofs"); + output(modsub.substr(1,1).c_str()); + + if(plusequalFlag == 1) { + output("+="); + } else { + output("+"); + } + + + if(modsub2[0]=='r') { + output("$"); + printCuobjdumpOperandlohi(modsub2); + } else { + modsub2 = modsub2.substr(2); + unsigned int addrValue; + std::stringstream hexStringConvert; + hexStringConvert << std::hex << modsub2; + hexStringConvert >> addrValue; + unsigned chunksize = 4; + if ( this->m_typeModifiers->getSize()>0 && + ( strcmp((m_typeModifiers->getListStart()->stringText), ".S16")==0 || + strcmp((m_typeModifiers->getListStart()->stringText), ".U16")==0 )) chunksize = 2; + if ( this->m_typeModifiers->getSize()>0 && + ( strcmp((m_typeModifiers->getListEnd()->stringText), ".U8")==0 || + strcmp((m_typeModifiers->getListEnd()->stringText), ".S8")==0 )) chunksize = 1; + addrValue = addrValue*chunksize; + char outputHex[10]; + sprintf(outputHex, "%x", addrValue); + std::stringstream outputhex; + outputhex << std::hex << addrValue; + output("0x"); + for(int i=4; i > outputhex.str().length(); i--) + { + output("0"); + } + output(outputhex.str().c_str()); + } + } else { + if(modsub[0]=='r') { + //Register + output("$"); + printCuobjdumpOperandlohi(modsub); + } else { + // Immediate offset + mod = mod.substr(2); + + unsigned int addrValue; + std::stringstream hexStringConvert; + hexStringConvert << std::hex << modsub; + hexStringConvert >> addrValue; + + if(localFlag == 0) + { + unsigned chunksize = 4; + if ( this->m_typeModifiers->getSize()>0 && + (strcmp((m_typeModifiers->getListStart()->stringText), ".S16")==0 || + strcmp((m_typeModifiers->getListStart()->stringText), ".U16")==0 )) chunksize = 2; + addrValue = addrValue*chunksize; + } + std::stringstream outputhex; + outputhex << std::hex << addrValue; + output("0x"); + for(int i=4; i > outputhex.str().length(); i--) { + output("0"); + } + output(outputhex.str().c_str()); + } + } + output("]"); + } else if(mod.substr(0,2) == "0x") { //immediate value + output("0x"); + std::string outputhex; + outputhex = mod.substr(2); + + for(int i=8; i > outputhex.length(); i--) { + output("0"); + } + output(outputhex.c_str()); + } else if(mod.substr(0,3) == "l0x") { //label + output(mod.c_str()); + } else {//variable name + output(mod.c_str()); + } +} + +void cuobjdumpInst::printCuobjdumpOperands() +{ + stringListPiece* currentPiece = m_operands->getListStart(); + std::string delimiter = ""; + for(int i=0; (i<m_operands->getSize())&&(currentPiece!=NULL); i++) + { + if((strcmp(m_base, "LOP.PASS_B")==0 || strcmp(m_base, "LOP.S.PASS_B")==0) && (i==1)) + { + currentPiece = currentPiece->nextString; + continue; + } + printCuobjdumpOperand(currentPiece, delimiter, m_base); + currentPiece = currentPiece->nextString; + delimiter = ","; + } +} + +char * int_default_mod () { return ".u32" ;} + + +std::string breaktarget; + +void cuobjdumpInst::printCuobjdumpPtxPlus(std::list<std::string> labelList, std::list<std::string> texList) +{ + stringListPiece* currentPiece; + + printCuobjdumpLabel(labelList); + + if(strcmp(m_base, "")==0) + { + } + else if(strcmp(m_base, ".entry")==0) + { + /*do nothing here*/ + } + else if(strcmp(m_base, "BAR.ARV.WAIT b0, 0xfff")==0) + { + printCuobjdumpPredicate(); + output("bar.sync 0x00000000;"); + } + else if(strcmp(m_base, "ADA")==0) + { + printCuobjdumpPredicate(); + output("add"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(int_default_mod()); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "BRA")==0) + { + printCuobjdumpPredicate(); + output("bra"); + printCuobjdumpBaseModifiers(); + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "CAL")==0) + { + printCuobjdumpPredicate(); + output("callp"); + printCuobjdumpBaseModifiers(); + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "COS")==0) + { + printCuobjdumpPredicate(); + //output("nop;"); + //output(" //cos"); + output("cos"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + } + else if(strcmp(m_base, "DADD")==0) + { + printCuobjdumpPredicate(); + output("add"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".ff64"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "DFMA")==0) + { + printCuobjdumpPredicate(); + output("mad"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".ff64"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "DMUL")==0) + { + printCuobjdumpPredicate(); + output("mul"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".ff64"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "EX2")==0) + { + printCuobjdumpPredicate(); + //output("nop;"); + //output(" //ex2"); + output("ex2"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "F2F")==0) + { + printCuobjdumpPredicate(); + + int absFlag = 0; + stringListPiece* currentPiece = m_baseModifiers->getListStart(); + + while(currentPiece != NULL) + { + if(strcmp(currentPiece->stringText, ".abs")==0) + { + output("abs"); + absFlag = 1; + break; + } + currentPiece = currentPiece->nextString; + } + if(absFlag == 0) + { + output("cvt"); + } + + printCuobjdumpBaseModifiers(); + + if(absFlag == 0) + { + printCuobjdumpTypeModifiers(); + } + else + { + output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right. + } + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "F2I")==0) + { + printCuobjdumpPredicate(); + output("cvt"); + printCuobjdumpBaseModifiers(); + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "FADD")==0 || strcmp(m_base, "FADD32I")==0) + { + printCuobjdumpPredicate(); + output("add"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "FADD32")==0) + { + printCuobjdumpPredicate(); + output("add.half"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "FMAD")==0 || strcmp(m_base, "FMAD32I")==0) + { + printCuobjdumpPredicate(); + output("mad"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "FMUL")==0 || strcmp(m_base, "FMUL32I")==0) + { + printCuobjdumpPredicate(); + output("mul"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "FMUL32")==0) + { + printCuobjdumpPredicate(); + output("mul.half"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "FSET")==0) + { + printCuobjdumpPredicate(); + output("set"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + { + output(".f32"); output(".f32"); + } + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "G2R")==0) + { + printCuobjdumpPredicate(); + output("mov"); + printCuobjdumpBaseModifiers(); + + if( m_typeModifiers->getSize() == 2 ) { + std::string type1, type2, type; + int type1Size, type2Size; + stringListPiece* currentPiece = m_typeModifiers->getListStart(); + type1 = currentPiece->stringText; + type2 = currentPiece->nextString->stringText; + + type1Size = atoi(type1.substr(2, type1.size()-2).c_str()); + type2Size = atoi(type2.substr(2, type2.size()-2).c_str()); + type = (type1Size < type2Size) ? type1 : type2; + + if( strcmp(type.c_str(), ".F16")==0 ) + output(".f16"); + else if( strcmp(type.c_str(), ".F32")==0 ) + output(".f32"); + else if( strcmp(type.c_str(), ".F64")==0 ) + output(".ff64"); + else if( strcmp(type.c_str(), ".S8")==0 ) + output(".s8"); + else if( strcmp(type.c_str(), ".S16")==0 ) + output(".s16"); + else if( strcmp(type.c_str(), ".S32")==0 ) + output(".s32"); + else if( strcmp(type.c_str(), ".S64")==0 ) + output(".bb64"); + else if( strcmp(type.c_str(), ".S128")==0 ) + output(".bb128"); + else if( strcmp(type.c_str(), ".U8")==0 ) + output(".u8"); + else if( strcmp(type.c_str(), ".U16")==0 ) + output(".u16"); + else if( strcmp(type.c_str(), ".U32")==0 ) + output(".u32"); + else if( strcmp(type.c_str(), ".U64")==0 ) + output(".bb64"); + else + output(type.c_str()); + } else if( m_typeModifiers->getSize() == 1 ) { + printCuobjdumpTypeModifiers(); + } else { + output("Error: unsupported number of type modifiers. "); + } + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "GLD")==0) + { + printCuobjdumpPredicate(); + //output("mov"); + output("ld.global"); + printCuobjdumpBaseModifiers(); + if (strcmp(m_typeModifiers->getListStart()->stringText, ".S128")==0) { + output(".v4.u32"); + } else if (strcmp(m_typeModifiers->getListStart()->stringText, ".S64")==0) { + output(".v2.u32"); + } else { + printCuobjdumpTypeModifiers(); + } + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "GST")==0) + { + printCuobjdumpPredicate(); + //output("mov"); + output("st.global"); + printCuobjdumpBaseModifiers(); + if (strcmp(m_typeModifiers->getListStart()->stringText, ".S128")==0) { + output(".v4.u32"); + } else if (strcmp(m_typeModifiers->getListStart()->stringText, ".S64")==0) { + output(".v2.u32"); + } else { + printCuobjdumpTypeModifiers(); + } + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "I2F")==0) + { + printCuobjdumpPredicate(); + output("cvt"); + printCuobjdumpBaseModifiers(); + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "I2I")==0) + { + printCuobjdumpPredicate(); + + int absFlag = 0; + int bextFlag = 0; + stringListPiece* currentPiece = m_baseModifiers->getListStart(); + + while(currentPiece != NULL) + { + if(strcmp(currentPiece->stringText, ".abs")==0) + { + output("abs"); + absFlag = 1; + break; + } + if(strcmp(currentPiece->stringText, ".bext")==0) + { + bextFlag = 1; + break; + } + currentPiece = currentPiece->nextString; + } + if(absFlag == 0) + { + output("cvt"); + } + + printCuobjdumpBaseModifiers(); + + currentPiece = this->m_typeModifiers->getListStart(); + unsigned int maxlength=16; + unsigned int currlength = 16; + bool issigned = false; + std::string tmpstr; + while(currentPiece != NULL) + { + tmpstr = currentPiece->stringText; + if(tmpstr[1] == 'S') + { + issigned = true; + } + if(tmpstr.substr(2,2) == "32") {currlength=32;} + if (currlength > maxlength) {maxlength = currlength;} + currentPiece = currentPiece->nextString; + } + + if(absFlag == 0) + { + if(bextFlag == 0) + printCuobjdumpTypeModifiers(); + else{ + output("."); + if (issigned) {output("s");} + else {output("u");} + std::stringstream tmp; + tmp << maxlength; + output(tmp.str().c_str()); + output("."); + if (issigned) {output("s");} + else {output("u");} + output("8"); + //output(".u16.u8"); + } + } + else + { + output(int_default_mod()); //TODO: setting default type modifier but I'm not sure if this is right. + } + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "IADD")==0) + { + printCuobjdumpPredicate(); + output("add"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(int_default_mod()); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "IADD32")==0 || strcmp(m_base, "IADD32I")==0 ) + { + printCuobjdumpPredicate(); + output("add.half"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".u32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "IMAD")==0 || strcmp(m_base, "IMAD32I")==0 || strcmp(m_base, "IMAD32")==0) + { + printCuobjdumpPredicate(); + output("mad.lo"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(int_default_mod()); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "IMAD.U24")==0) + { + printCuobjdumpPredicate(); + output("mad24.lo"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".u32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "IMAD.S24")==0) + { + printCuobjdumpPredicate(); + output("mad24.lo"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".s32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "IMUL32I.U24.U24")==0 || + strcmp(m_base, "IMUL32I.S24.S24")==0 || + strcmp(m_base, "IMUL32.U24.U24")==0 || + strcmp(m_base, "IMUL32.S24.S24")==0 || + strcmp(m_base, "IMUL.U24.U24")==0 || + strcmp(m_base, "IMUL.S24.S24")==0) + { + printCuobjdumpPredicate(); + output("mul24.lo"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + { + output(".u32"); //TODO: setting default type modifier but I'm not sure if this is right. + } + else if(m_typeModifiers->getSize() == 2) + { + std::string type1, type2, type; + int type1Size, type2Size; + stringListPiece* currentPiece = m_typeModifiers->getListStart(); + char tempString[5]; + + type1 = currentPiece->stringText; + type2 = currentPiece->nextString->stringText; + + type1Size = atoi(type1.substr(2, type1.size()-2).c_str()); + type2Size = atoi(type2.substr(2, type2.size()-2).c_str()); + type = (type1Size > type2Size) ? type1 : type2; + strcpy(tempString, type.c_str()); + /*if(type1Size==16 && type2Size==16) + output(".lo");*/ + if(tempString[1] >= 'A' && tempString[1] <= 'Z') + tempString[1] += 32; + output(tempString); + } + else + { + printCuobjdumpTypeModifiers(); + } + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "IMUL")==0 || strcmp(m_base, "IMUL32I")==0) + { + printCuobjdumpPredicate(); + output("mul.lo"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + { + output(int_default_mod()); //TODO: setting default type modifier but I'm not sure if this is right. + } + else if(m_typeModifiers->getSize() == 2) + { + std::string type1, type2, type; + int type1Size, type2Size; + stringListPiece* currentPiece = m_typeModifiers->getListStart(); + char tempString[5]; + + type1 = currentPiece->stringText; + type2 = currentPiece->nextString->stringText; + + type1Size = atoi(type1.substr(2, type1.size()-2).c_str()); + type2Size = atoi(type2.substr(2, type2.size()-2).c_str()); + type = (type1Size > type2Size) ? type1 : type2; + strcpy(tempString, type.c_str()); + /*if(type1Size==16 && type2Size==16) + output(".lo");*/ + if(tempString[1] >= 'A' && tempString[1] <= 'Z') + tempString[1] += 32; + output(tempString); + } + else + { + printCuobjdumpTypeModifiers(); + } + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "IMUL32")==0) + { + printCuobjdumpPredicate(); + output("mul.half.lo"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 2) + { + std::string type1, type2, type; + int type1Size, type2Size; + stringListPiece* currentPiece = m_typeModifiers->getListStart(); + char tempString[5]; + + type1 = currentPiece->stringText; + type2 = currentPiece->nextString->stringText; + + type1Size = atoi(type1.substr(2, type1.size()-2).c_str()); + type2Size = atoi(type2.substr(2, type2.size()-2).c_str()); + type = (type1Size > type2Size) ? type1 : type2; + strcpy(tempString, type.c_str()); + /*if(type1Size==16 && type2Size==16) + output(".lo");*/ + if(tempString[1] >= 'A' && tempString[1] <= 'Z') + tempString[1] += 32; + output(tempString); + + } + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "IMUL32.S24.S24")==0) + { + printCuobjdumpPredicate(); + output("mul24.half.lo"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + { + output(".u32"); //TODO: setting default type modifier but I'm not sure if this is right. + } + else if(m_typeModifiers->getSize() == 2) + { + std::string type1, type2, type; + int type1Size, type2Size; + stringListPiece* currentPiece = m_typeModifiers->getListStart(); + char tempString[5]; + + type1 = currentPiece->stringText; + type2 = currentPiece->nextString->stringText; + + type1Size = atoi(type1.substr(2, type1.size()-2).c_str()); + type2Size = atoi(type2.substr(2, type2.size()-2).c_str()); + type = (type1Size > type2Size) ? type1 : type2; + strcpy(tempString, type.c_str()); + /*if(type1Size==16 && type2Size==16) + output(".lo");*/ + if(tempString[1] >= 'A' && tempString[1] <= 'Z') + tempString[1] += 32; + output(tempString); + + } + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "ISET")==0) + { + printCuobjdumpPredicate(); + output("set"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".u32.u32"); //TODO: setting default type modifier but I'm not sure if this is right. + else if(m_typeModifiers->getSize() == 1) + { + printCuobjdumpTypeModifiers(); printCuobjdumpTypeModifiers(); + } + else + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "LG2")==0) + { + printCuobjdumpPredicate(); + output("lg2"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "LST")==0) + { + printCuobjdumpPredicate(); + output("mov"); + printCuobjdumpBaseModifiers(); + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "MVC")==0) { + printCuobjdumpPredicate(); + output("cvt"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".u32.u32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + output(";"); + } + + else if(strcmp(m_base, "MOV")==0 || strcmp(m_base, "MVI")==0 || strcmp(m_base, "MVC")==0) + { + printCuobjdumpPredicate(); + output("mov"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(int_default_mod()); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "MOV32")==0) + { + printCuobjdumpPredicate(); + output("mov.half"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".u32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "NOP")==0 || strcmp(m_base, "SSY")==0) + { + printCuobjdumpPredicate(); + output("nop"); + printCuobjdumpBaseModifiers(); + output(";"); + } + else if(strcmp(m_base, "LLD")==0) + { + printCuobjdumpPredicate(); + output("mov"); + printCuobjdumpBaseModifiers(); + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "LOP.AND")==0 || strcmp(m_base, "LOP.S.AND")==0) + { + printCuobjdumpPredicate(); + output("and"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".b32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "LOP.OR")==0 || strcmp(m_base, "LOP.S.OR")==0) + { + printCuobjdumpPredicate(); + output("or"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".b32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "LOP.PASS_B")==0 || strcmp(m_base, "LOP.S.PASS_B")==0) + { + printCuobjdumpPredicate(); + output("not"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".b32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "LOP.XOR")==0 || strcmp(m_base, "LOP.S.XOR")==0) + { + printCuobjdumpPredicate(); + output("xor"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".b32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "R2A")==0) + { + printCuobjdumpPredicate(); + output("shl"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".b32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + if(m_operands->getSize() == 2) + output(", 0x0"); + output(";"); + } + else if(strcmp(m_base, "R2G.U16.U8")==0){ + /* + * This code handles a cuobjdump bug that causes the wrong register number to be printed + */ + printCuobjdumpPredicate(); + output("mov.u8"); + stringListPiece* currentPiece = m_operands->getListStart(); + std::string delimiter = ""; + printCuobjdumpOperand(currentPiece, delimiter, m_base); + output(", "); + std::string curr = currentPiece->nextString->stringText; + curr = curr.substr(1); + int regnum; + std::istringstream(curr)>>regnum; + output("$r"); + std::stringstream finalregnum; + finalregnum << (regnum/2); + output(finalregnum.str().c_str()); + output( regnum%2==0? ".lo": ".hi"); + output(";"); + } + else if(strcmp(m_base, "R2G")==0) + { + printCuobjdumpPredicate(); + output("mov"); + printCuobjdumpBaseModifiers(); + + if( m_typeModifiers->getSize() == 2 ) { + std::string type1, type2, type; + int type1Size, type2Size; + stringListPiece* currentPiece = m_typeModifiers->getListStart(); + type1 = currentPiece->stringText; + type2 = currentPiece->nextString->stringText; + + type1Size = atoi(type1.substr(2, type1.size()-2).c_str()); + type2Size = atoi(type2.substr(2, type2.size()-2).c_str()); + type = (type1Size < type2Size) ? type1 : type2; + + if( strcmp(type.c_str(), ".F16")==0 ) + output(".f16"); + else if( strcmp(type.c_str(), ".F32")==0 ) + output(".f32"); + else if( strcmp(type.c_str(), ".F64")==0 ) + output(".ff64"); + else if( strcmp(type.c_str(), ".S8")==0 ) + output(".s8"); + else if( strcmp(type.c_str(), ".S16")==0 ) + output(".s16"); + else if( strcmp(type.c_str(), ".S32")==0 ) + output(".s32"); + else if( strcmp(type.c_str(), ".S64")==0 ) + output(".bb64"); + else if( strcmp(type.c_str(), ".S128")==0 ) + output(".bb128"); + else if( strcmp(type.c_str(), ".U8")==0 ) + output(".u8"); + else if( strcmp(type.c_str(), ".U16")==0 ) + output(".u16"); + else if( strcmp(type.c_str(), ".U32")==0 ) + output(".u32"); + else if( strcmp(type.c_str(), ".U64")==0 ) + output(".bb64"); + else + output(type.c_str()); + } else if( m_typeModifiers->getSize() == 1 ) { + printCuobjdumpTypeModifiers(); + } else { + output("Error: unsupported number of type modifiers. "); + } + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "RCP")==0) + { + printCuobjdumpPredicate(); + output("rcp"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "RCP32")==0) + { + printCuobjdumpPredicate(); + output("rcp.half"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".f32"); + else + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "RET")==0) + { + printCuobjdumpPredicate(); + output("retp;"); + } + else if(strcmp(m_base, "RRO")==0) + { + output("nop; //"); + printCuobjdumpPredicate(); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "RSQ")==0) + { + printCuobjdumpPredicate(); + output("rsqrt"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "SHL")==0) + { + printCuobjdumpPredicate(); + output("shl"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(int_default_mod()); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "SHR")==0) + { + printCuobjdumpPredicate(); + output("shr"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(int_default_mod()); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "SIN")==0) + { + printCuobjdumpPredicate(); + output("sin"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + } + else if(strcmp(m_base, "S2R")==0) + { + printCuobjdumpPredicate(); + output("cvt"); + printCuobjdumpBaseModifiers(); + + if(m_typeModifiers->getSize() == 0) + output(".u32"); + else + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + } + else if(strcmp(m_base, "LD")==0) + { + printCuobjdumpPredicate(); + output("mov"); + printCuobjdumpBaseModifiers(); + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "STS")==0) + { + printCuobjdumpPredicate(); + output("mov"); + printCuobjdumpBaseModifiers(); + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "BAR")==0) + { + printCuobjdumpPredicate(); + output("bar.sync 0;"); + } + else if(strcmp(m_base, "LDS")==0) + { + // If there is not global address space that includes shared memory, then fix this. + // Same for STS + printCuobjdumpPredicate(); + output("mov"); + printCuobjdumpBaseModifiers(); + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "ST")==0) + { + printCuobjdumpPredicate(); + output("mov"); + printCuobjdumpBaseModifiers(); + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + output(";"); + } else if(strcmp(m_base, "IMIN")==0) { + printCuobjdumpPredicate(); + output("min"); + printCuobjdumpBaseModifiers(); + if(m_typeModifiers->getSize() == 0) + output(".s32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + output(";"); + } else if(strcmp(m_base, "IMAX")==0) { + printCuobjdumpPredicate(); + output("max"); + printCuobjdumpBaseModifiers(); + if(m_typeModifiers->getSize() == 0) + output(".s32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + output(";"); + } else if(strcmp(m_base, "FMIN")==0) { + printCuobjdumpPredicate(); + output("min"); + printCuobjdumpBaseModifiers(); + if(m_typeModifiers->getSize() == 0) + output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "FMAX")==0) { + printCuobjdumpPredicate(); + output("max"); + printCuobjdumpBaseModifiers(); + if(m_typeModifiers->getSize() == 0) + output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + output(";"); + } + else if(strcmp(m_base, "A2R")==0) { + printCuobjdumpPredicate(); + output("mov"); + printCuobjdumpBaseModifiers(); + if(m_typeModifiers->getSize() == 0) + output(int_default_mod()); //TODO: setting default type modifier but I'm not sure if this is right. + else + printCuobjdumpTypeModifiers(); + printCuobjdumpOperands(); + output(";"); + } + else if((strcmp(m_base, "TEX")==0) || + (strcmp(m_base, "TEX32")==0)) { + printCuobjdumpPredicate(); + output("tex.1d.v4.f32.s32 "); + stringListPiece *currPiece; + stringListPiece reg; + std::string addrReg, tex_id; + currPiece = m_operands->getListStart(); + output("{"); + printCuobjdumpOperand(currPiece, "", ""); + output(",_,_,_} , "); + reg = *currPiece; + currPiece = currPiece->nextString; + tex_id = currPiece->stringText; + unsigned int tex_id_int; + std::stringstream ss; + ss << std::hex << tex_id; + ss >> tex_id_int; + std::list<std::string>::iterator texIter = texList.begin(); + for (int i=0; i<tex_id_int; i++) { + assert (texIter!=texList.end()); + texIter++; + } + output((*texIter).c_str()); + output(",{"); + printCuobjdumpOperand(®, "", ""); + output(",_,_,_};"); + } + else if(strcmp(m_base, "EXIT")==0) { + printCuobjdumpPredicate(); + output("exit"); + output(";"); + } + else if(strcmp(m_base, "GRED")==0) { + printCuobjdumpPredicate(); + // ptx instruction atom can be used to perform reduction using destination register '_' + output("atom.global"); + printCuobjdumpBaseModifiers(); + printCuobjdumpTypeModifiers(); + output(" _, "); + printCuobjdumpOperands(); + output(";"); + } else if(strcmp(m_base, "GATOM")==0) { + printCuobjdumpPredicate(); + output("atom.global"); + printCuobjdumpBaseModifiers(); + printCuobjdumpTypeModifiers(); + + printCuobjdumpOperands(); + output(";"); + }else if(strcmp(m_base, "PBK")==0) { + //PKB specifies the target of later BRK (break) instructions + // Here we convert it to nop and store the target + output("nop;"); + breaktarget = m_operands->getListStart()->stringText; + } else if(strcmp(m_base, "BRK")==0) { + printCuobjdumpPredicate(); + /* + * Convert it **at compile time** to a branch to the break target saved earlier + * Correct operation not guaranteed + */ + + output("bra "); + output(breaktarget.c_str()); + output(";"); + } else if(strcmp(m_base, "C2R")==0) { + printCuobjdumpPredicate(); + output("mov.u32"); + printCuobjdumpOperands(); + output(";"); + } else if(strcmp(m_base, "R2C")==0) { + printCuobjdumpPredicate(); + output("mov.pred"); + printCuobjdumpOperands(); + output(";"); + } + else + { + printf("Unknown Instruction: "); + printf(m_base); + printf("\n"); + output("Unknown Instruction: "); + output(m_base); + assert(0); + } + +} + |