diff options
Diffstat (limited to 'decuda_to_ptxplus/decudaInst.cc')
| -rw-r--r-- | decuda_to_ptxplus/decudaInst.cc | 1523 |
1 files changed, 0 insertions, 1523 deletions
diff --git a/decuda_to_ptxplus/decudaInst.cc b/decuda_to_ptxplus/decudaInst.cc deleted file mode 100644 index 4f1b9de..0000000 --- a/decuda_to_ptxplus/decudaInst.cc +++ /dev/null @@ -1,1523 +0,0 @@ -// Copyright (c) 2009-2011, Jimmy Kwa, -// The University of British Columbia -// All rights reserved. -// -// Redistribution and use in source and binary forms, with or without -// modification, are permitted provided that the following conditions are met: -// -// Redistributions of source code must retain the above copyright notice, this -// list of conditions and the following disclaimer. -// Redistributions in binary form must reproduce the above copyright notice, this -// list of conditions and the following disclaimer in the documentation and/or -// other materials provided with the distribution. -// Neither the name of The University of British Columbia nor the names of its -// contributors may be used to endorse or promote products derived from this -// software without specific prior written permission. -// -// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND -// ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED -// WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE -// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE -// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL -// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR -// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER -// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, -// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE -// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. - - -#include "decudaInst.h" -#include <iostream> -#include <string> -#include <cstring> -#include <stdlib.h> -#include <stdio.h> - -extern void output(const char * text); - -//Constructor -decudaInst::decudaInst() -{ - //initilize everything to empty - m_label = ""; - m_predicate = new stringList(); - m_base = ""; - m_baseModifiers = new stringList(); - m_typeModifiers = new stringList(); - m_operands = new stringList(); - m_predicateModifiers = new stringList(); - - m_nextDecudaInst = NULL; - - // Set operations per cycle to 8 by default (no penalty) - m_opPerCycle = 8; -} - -//retreive instruction mnemonic -const char* decudaInst::getBase() -{ - return m_base; -} - -stringList* decudaInst::getOperands() -{ - return m_operands; -} - -stringList* decudaInst::getBaseModifiers() -{ - return m_baseModifiers; -} - -stringList* decudaInst::getTypeModifiers() -{ - return m_typeModifiers; -} - -//get next instruction in linked list -//direction is m_listStart to m_listEnd -decudaInst* decudaInst::getNextDecudaInst() -{ - return m_nextDecudaInst; -} - -void decudaInst::setBase(const char* setBaseValue) -{ - m_base = setBaseValue; -} - -void decudaInst::addBaseModifier(const char* addBaseMod) -{ - stringListPiece* tempPiece = new stringListPiece; - tempPiece->stringText = addBaseMod; - - m_baseModifiers->add(tempPiece); -} - -void decudaInst::addTypeModifier(const char* addTypeMod) -{ - stringListPiece* tempPiece = new stringListPiece; - tempPiece->stringText = addTypeMod; - - m_typeModifiers->add(tempPiece); -} - -void decudaInst::addOperand(const char* addOp) -{ - stringListPiece* tempPiece = new stringListPiece; - tempPiece->stringText = addOp; - - m_operands->add(tempPiece); -} - -void decudaInst::setPredicate(const char* setPredicateValue) -{ - stringListPiece* tempPiece = new stringListPiece; - tempPiece->stringText = setPredicateValue; - - m_predicate->add(tempPiece); -} - -void decudaInst::addPredicateModifier(const char* addPredicateMod) -{ - stringListPiece* tempPiece = new stringListPiece; - tempPiece->stringText = addPredicateMod; - - m_predicateModifiers->add(tempPiece); -} - -void decudaInst::setLabel(const char* setLabelValue) -{ - m_label = setLabelValue; -} - -//set next instruction in linked list -//direction is m_listStart to m_listEnd -void decudaInst::setNextDecudaInst(decudaInst* setDecudaInstValue) -{ - m_nextDecudaInst = setDecudaInstValue; -} - -// returns true if current instruction is start of an entry (i.e. '{') -bool decudaInst::isEntryStart() -{ - return (strcmp(m_base, "{")==0); -} - -//print out .version and .target header lines -bool decudaInst::printHeaderInst() -{ - if(strcmp(m_base, ".version")==0) - { - output(m_base); output(" "); - - stringListPiece* currentPiece = m_operands->getListStart(); - output(currentPiece->stringText); - currentPiece = currentPiece->nextString; - - if(currentPiece!=NULL) - { - output("."); output(currentPiece->stringText); - } - output("+"); - output("\n"); - } - else if(strcmp(m_base, ".target")==0) - { - output(m_base); output(" "); - - stringListPiece* currentPiece = m_operands->getListStart(); - output(currentPiece->stringText); - currentPiece = currentPiece->nextString; - - while(currentPiece!=NULL) - { - output(", "); output(currentPiece->stringText); - currentPiece = currentPiece->nextString; - } - output("\n"); - } - else if(strcmp(m_base, ".tex")==0) - { - output(m_base); output(" "); - - stringListPiece* currentPiece; - - - currentPiece = m_baseModifiers->getListStart(); - output(currentPiece->stringText); output(" "); - currentPiece = currentPiece->nextString; - - while(currentPiece!=NULL) - { - output(" "); output(currentPiece->stringText); - currentPiece = currentPiece->nextString; - } - - currentPiece = m_operands->getListStart(); - output(currentPiece->stringText); - currentPiece = currentPiece->nextString; - - while(currentPiece!=NULL) - { - output(" "); output(currentPiece->stringText); - currentPiece = currentPiece->nextString; - } - output(";\n"); - } - else - { - return false; - } - return true; -} - -//print out parameters -bool decudaInst::printHeaderInst2() -{ - if(strcmp(m_base, ".param")==0) - { - output(m_base); - - stringListPiece* currentPiece = m_typeModifiers->getListStart(); - - for(int i=0; (i<m_typeModifiers->getSize())&&(currentPiece!=NULL); i++) - { - output(" "); output(currentPiece->stringText); - - currentPiece = currentPiece->nextString; - } - - currentPiece = m_operands->getListStart(); - for(int i=0; (i<m_operands->getSize())&&(currentPiece!=NULL); i++) - { - output(" "); output(currentPiece->stringText); - - currentPiece = currentPiece->nextString; - } - } - else - { - return false; - } - return true; -} -//print out the Decuda instruction -void decudaInst::printDecudaInst() -{ - /*TODO: print label here*/ - /*TODO: print predicate here*/ - /*std::cout << "Instruction Base: " << m_base << "\n"; - - std::cout << "Instruction Modifiers: "; - m_baseModifiers->printStringList(); - std::cout << "\n"; - - std::cout << "Operand types: "; - m_typeModifiers->printStringList(); - std::cout << "\n"; - - std::cout << "Operands: "; - m_operands->printStringList(); - std::cout << "\n\n";*/ - - std::cout << m_base << " "; - m_baseModifiers->printStringList(); - std::cout << " "; - m_typeModifiers->printStringList(); - std::cout << " "; - m_operands->printStringList(); - std::cout << "\n"; - - -} - - -// Just prints the base and operands -void decudaInst::printHeaderPtx() -{ - output(m_base); output(" "); - - stringListPiece* currentPiece; - - currentPiece = m_baseModifiers->getListStart(); - for(int i=0; (i<m_baseModifiers->getSize())&&(currentPiece!=NULL); i++) - { - output(" "); output(currentPiece->stringText); - currentPiece = currentPiece->nextString; - } - - currentPiece = m_operands->getListStart(); - for(int i=0; (i<m_operands->getSize())&&(currentPiece!=NULL); i++) - { - output(" "); output(currentPiece->stringText); - currentPiece = currentPiece->nextString; - } -} - - - -// Print unmodified instruction -void decudaInst::printDefaultPtx() -{ - printLabel(); - printPredicate(); - - output(m_base); - - printBaseModifiers(); - printTypeModifiers(); - printOperands(); - - output(";"); -} - - -// Print unmodified base modifiers, operands, labels and predicates -void decudaInst::printLabel() -{ - if(m_label != "") { - output(m_label); - output(": "); - } -} - -void decudaInst::printPredicate() -{ - stringListPiece* currentPiece = m_predicate->getListStart(); - if(currentPiece!=NULL) - { - output(currentPiece->stringText); - - stringListPiece* currentPiece2 = m_predicateModifiers->getListStart(); - for(int i=0; (i<m_predicateModifiers->getSize())&&(currentPiece2!=NULL); i++) - { - output(currentPiece2->stringText); - - currentPiece2 = currentPiece2->nextString; - } - - output(" "); - } -} - -void decudaInst::printBaseModifiers() -{ - stringListPiece* currentPiece = m_baseModifiers->getListStart(); - for(int i=0; (i<m_baseModifiers->getSize())&&(currentPiece!=NULL); i++) - { - output(currentPiece->stringText); - - currentPiece = currentPiece->nextString; - } -} - -void decudaInst::printTypeModifiers() -{ - stringListPiece* currentPiece = m_typeModifiers->getListStart(); - for(int i=0; (i<m_typeModifiers->getSize())&&(currentPiece!=NULL); i++) - { - output(currentPiece->stringText); - - currentPiece = currentPiece->nextString; - } -} - -void decudaInst::printOperands() -{ - stringListPiece* currentPiece = m_operands->getListStart(); - const char* operandDelimiter = ""; - for(int i=0; (i<m_operands->getSize())&&(currentPiece!=NULL); i++) - { - output(operandDelimiter); output(" "); output(currentPiece->stringText); - operandDelimiter = ","; - - currentPiece = currentPiece->nextString; - } -} - -//This is where the conversion to new PTX takes place -void decudaInst::printNewPtx() -{ - // - // Common modifications that apply to all instructions - // - stringListPiece* currentPiece; - int vectorFlag[4] = {0,0,0,0}; //0=16/32type, 1=bb64/ff64 type, 2=bb128 type - - // Replace '%clock' with '%halfclock' - currentPiece = m_operands->getListStart(); - for(int i=0; (i<m_operands->getSize())&&(currentPiece!=NULL); i++) - { - const char* modString = currentPiece->stringText; - if( strcmp(modString, "%%clock")==0 ) { - const char* newText = "%%halfclock"; - currentPiece->stringText = newText; - } - currentPiece = currentPiece->nextString; - } - - // Remove .join from base modifier list - currentPiece = m_baseModifiers->getListStart(); - for(int i=0; (i<m_baseModifiers->getSize())&&(currentPiece!=NULL); i++) - { - const char* modString = currentPiece->stringText; - if( strcmp(modString, ".join")==0 ) { - m_baseModifiers->remove(i); - } - currentPiece = currentPiece->nextString; - } - - // Change .end to .exit from base modified list - currentPiece = m_baseModifiers->getListStart(); - for(int i=0; (i<m_baseModifiers->getSize())&&(currentPiece!=NULL); i++) - { - const char* modString = currentPiece->stringText; - if( strcmp(modString, ".end")==0 ) { - const char* newText = ".exit"; - currentPiece->stringText = newText; - } - currentPiece = currentPiece->nextString; - } - - // Change .b64 to .bb64 from type modifier list - // Change .b128 to .bb128 from type modifier list - // Change .f64 to .ff64 from type modifier list - currentPiece = m_typeModifiers->getListStart(); - for(int i=0; (i<m_typeModifiers->getSize())&&(currentPiece!=NULL); i++) - { - const char* modString = currentPiece->stringText; - if( strcmp(modString, ".b64")==0 ) { - const char* newText = ".bb64"; - currentPiece->stringText = newText; - vectorFlag[i] = 1; - } - else if( strcmp(modString, ".b128")==0 ) { - const char* newText = ".bb128"; - currentPiece->stringText = newText; - vectorFlag[i] = 2; - } - else if( strcmp(modString, ".f64")==0 ) { - const char* newText = ".ff64"; - currentPiece->stringText = newText; - vectorFlag[i] = 1; - } - currentPiece = currentPiece->nextString; - } - - /*decuda bug workaround. - cvt.abs should drop the first type modifier. - cvt.rn.f32.f64 needs to have type modifiers reversed to cvt.ff64.f32*/ - int absFound = 0; - if(strcmp(m_base, "cvt")==0 && m_typeModifiers->getSize() == 2) - { - stringListPiece* currentPieceCvt = m_baseModifiers->getListStart(); - for(int i=0; (i<m_baseModifiers->getSize())&&(currentPieceCvt!=NULL); i++) - { - const char* modStringCvt = currentPieceCvt->stringText; - if( strcmp(modStringCvt, ".abs")==0 ) { - vectorFlag[0] = vectorFlag[1]; - absFound = 1; - break; - } - currentPieceCvt = currentPieceCvt->nextString; - } - if(absFound == 0 && vectorFlag[1] == 1) - { - vectorFlag[0] = 1; - vectorFlag[1] = 0; - const char* tempCharPtr = m_typeModifiers->getListStart()->stringText; - m_typeModifiers->getListStart()->stringText = m_typeModifiers->getListStart()->nextString->stringText; - m_typeModifiers->getListStart()->nextString->stringText = tempCharPtr; - } - } - - /*decuda bug workaround, cvt.rz.f64 is really cvt.rz.f32.f64*/ - if(strcmp(m_base, "cvt")==0 && m_typeModifiers->getSize() == 1 && vectorFlag[0]==1) - { - vectorFlag[0] = 0; - vectorFlag[1] = 1; - addTypeModifier(m_typeModifiers->getListStart()->stringText); - const char* newText = ".f32"; - m_typeModifiers->getListStart()->stringText = newText; - } - - /*expand vector operands eg. $r0 -> {$r0, $r1}*/ - if((vectorFlag[0] != 0) || (vectorFlag[1] != 0) || (vectorFlag[2] != 0) || (vectorFlag[3] != 0) ) - { - currentPiece = m_operands->getListStart(); - for(int i=0; (i<m_operands->getSize())&&(currentPiece!=NULL); i++) - { - char *newText = new char[40]; - char *regNumString; - int regNumInt; - - const char* modString = currentPiece->stringText; - - if(strcmp(m_base, "set")==0 || strcmp(m_base, "cvt")==0 || - strcmp(m_base, "set?68?")==0 || strcmp(m_base, "set?65?")==0 || - strcmp(m_base, "set?67?")==0 || strcmp(m_base, "set?13?")==0) - { - if( modString[0] == '$' && modString[1] == 'r' ) { - strcpy(newText, modString); - strtok (newText, "r"); - regNumString = strtok (NULL, "r"); - regNumInt = atoi(regNumString); - if(vectorFlag[i] ==0) - strcpy(newText, modString); - if(vectorFlag[i] ==1) - snprintf(newText,40,"{$r%u,$r%u}", regNumInt+0, regNumInt+1); - if(vectorFlag[i] ==2) - snprintf(newText,40,"{$r%u,$r%u,$r%u,$r%u}", regNumInt+0, regNumInt+1, regNumInt+2, regNumInt+3); - - currentPiece->stringText = newText; - } else if( modString[0] == '-' && modString[1] == '$' && modString[2] == 'r' ) { - strcpy(newText, modString); - strtok (newText, "r"); - regNumString = strtok (NULL, "r"); - regNumInt = atoi(regNumString); - if(vectorFlag[i] ==0) - strcpy(newText, modString); - else if(vectorFlag[i] ==1) - snprintf(newText,40,"-{$r%u,$r%u}", regNumInt+0, regNumInt+1); - else if(vectorFlag[i] ==2) - snprintf(newText,40,"-{$r%u,$r%u,$r%u,$r%u}", regNumInt+0, regNumInt+1, regNumInt+2, regNumInt+3); - - currentPiece->stringText = newText; - } - } - else - { - if( modString[0] == '$' && modString[1] == 'r' ) { - strcpy(newText, modString); - strtok (newText, "r"); - regNumString = strtok (NULL, "r"); - regNumInt = atoi(regNumString); - if(vectorFlag[0] ==0) - strcpy(newText, modString); - else if(vectorFlag[0] ==1) - snprintf(newText,40,"{$r%u,$r%u}", regNumInt+0, regNumInt+1); - else if(vectorFlag[0] ==2) - snprintf(newText,40,"{$r%u,$r%u,$r%u,$r%u}", regNumInt+0, regNumInt+1, regNumInt+2, regNumInt+3); - - currentPiece->stringText = newText; - } else if( modString[0] == '-' && modString[1] == '$' && modString[2] == 'r' ) { - strcpy(newText, modString); - strtok (newText, "r"); - regNumString = strtok (NULL, "r"); - regNumInt = atoi(regNumString); - if(vectorFlag[0] ==0) - strcpy(newText, modString); - else if(vectorFlag[0] ==1) - snprintf(newText,40,"-{$r%u,$r%u}", regNumInt+0, regNumInt+1); - else if(vectorFlag[0] ==2) - snprintf(newText,40,"-{$r%u,$r%u,$r%u,$r%u}", regNumInt+0, regNumInt+1, regNumInt+2, regNumInt+3); - - currentPiece->stringText = newText; - } - } - currentPiece = currentPiece->nextString; - } - } - - // - // Instruction specific modifications - // - if(strcmp(m_base, "")==0) - { - } - else if(strcmp(m_base, ".entry")==0) - { - /*do nothing here*/ - } - else if(strcmp(m_base, ".lmem")==0) - { - } - else if(strcmp(m_base, ".smem")==0) - { - } - else if(strcmp(m_base, ".reg")==0) - { - } - else if(strcmp(m_base, ".bar")==0) - { - } - - - else if(strcmp(m_base, "cvt")==0) - { - - int cvt_inst_type = 0; //0==cvt, 1==abs - int cvt_neg_mod_flag = 0; //0=off, 1=on - - // Check the actual base instruction - stringListPiece* currentPiece = m_baseModifiers->getListStart(); - for(int i=0; (i<m_baseModifiers->getSize())&&(currentPiece!=NULL); i++) - { - const char* modString = currentPiece->stringText; - if( strcmp(modString, ".abs")==0 ) { - cvt_inst_type = 1; - } - if( strcmp(modString, ".neg")==0 ) { - cvt_neg_mod_flag = 1; - } - currentPiece = currentPiece->nextString; - } - - if( cvt_inst_type == 1) { - // abs instruction - printLabel(); - printPredicate(); - - output("abs"); - - if(m_typeModifiers->getSize() == 1) - { - const char* type = m_typeModifiers->getListStart()->stringText; - output(type); - } - else - { - const char* type = m_typeModifiers->getListStart()->nextString->stringText; - output(type); - } - - printOperands(); - - output(";"); - - } else { - // cvt instruction - printLabel(); - printPredicate(); - - output("cvt"); - - int typeModifiers = m_typeModifiers->getSize(); - - if(typeModifiers == 2) { - - stringListPiece* currentPiece = m_baseModifiers->getListStart(); - - if(cvt_neg_mod_flag != 0) - currentPiece = m_baseModifiers->getListStart()->nextString; - - const char* dstType = m_typeModifiers->getListStart()->stringText; - const char* srcType = m_typeModifiers->getListStart()->nextString->stringText; - - for(int i=0; (i<m_baseModifiers->getSize())&&(currentPiece!=NULL); i++) - { - const char* modString = currentPiece->stringText; - if( - (strcmp(modString, ".rn")==0) || - (strcmp(modString, ".rm")==0) || - (strcmp(modString, ".rp")==0) || - (strcmp(modString, ".rz")==0) - ) - { - if((strcmp(dstType, ".f32")==0) && (strcmp(srcType, ".ff64")==0)) - output(modString); - if((strcmp(dstType, ".f16")==0) && (strcmp(srcType, ".ff64")==0)) - output(modString); - if((strcmp(dstType, ".f16")==0) && (strcmp(srcType, ".f32")==0)) - output(modString); - if((strcmp(dstType, ".f16")==0) && (strcmp(srcType, ".u8")==0)) - output(modString); - if((strcmp(dstType, ".f16")==0) && (strcmp(srcType, ".u16")==0)) - output(modString); - if((strcmp(dstType, ".f16")==0) && (strcmp(srcType, ".u32")==0)) - output(modString); - if((strcmp(dstType, ".f16")==0) && (strcmp(srcType, ".u64")==0)) - output(modString); - if((strcmp(dstType, ".f16")==0) && (strcmp(srcType, ".s8")==0)) - output(modString); - if((strcmp(dstType, ".f16")==0) && (strcmp(srcType, ".s16")==0)) - output(modString); - if((strcmp(dstType, ".f16")==0) && (strcmp(srcType, ".s32")==0)) - output(modString); - if((strcmp(dstType, ".f16")==0) && (strcmp(srcType, ".s64")==0)) - output(modString); - if((strcmp(dstType, ".f32")==0) && (strcmp(srcType, ".u8")==0)) - output(modString); - if((strcmp(dstType, ".f32")==0) && (strcmp(srcType, ".u16")==0)) - output(modString); - if((strcmp(dstType, ".f32")==0) && (strcmp(srcType, ".u32")==0)) - output(modString); - if((strcmp(dstType, ".f32")==0) && (strcmp(srcType, ".u64")==0)) - output(modString); - if((strcmp(dstType, ".f32")==0) && (strcmp(srcType, ".s8")==0)) - output(modString); - if((strcmp(dstType, ".f32")==0) && (strcmp(srcType, ".s16")==0)) - output(modString); - if((strcmp(dstType, ".f32")==0) && (strcmp(srcType, ".s32")==0)) - output(modString); - if((strcmp(dstType, ".f32")==0) && (strcmp(srcType, ".s64")==0)) - output(modString); - if((strcmp(dstType, ".ff64")==0) && (strcmp(srcType, ".u8")==0)) - output(modString); - if((strcmp(dstType, ".ff64")==0) && (strcmp(srcType, ".u16")==0)) - output(modString); - if((strcmp(dstType, ".ff64")==0) && (strcmp(srcType, ".u32")==0)) - output(modString); - if((strcmp(dstType, ".ff64")==0) && (strcmp(srcType, ".u64")==0)) - output(modString); - if((strcmp(dstType, ".ff64")==0) && (strcmp(srcType, ".s8")==0)) - output(modString); - if((strcmp(dstType, ".ff64")==0) && (strcmp(srcType, ".s16")==0)) - output(modString); - if((strcmp(dstType, ".ff64")==0) && (strcmp(srcType, ".s32")==0)) - output(modString); - if((strcmp(dstType, ".ff64")==0) && (strcmp(srcType, ".s64")==0)) - output(modString); - if((strcmp(dstType, ".u8")==0) && (strcmp(srcType, ".f16")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".u16")==0) && (strcmp(srcType, ".f16")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".u32")==0) && (strcmp(srcType, ".f16")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".u64")==0) && (strcmp(srcType, ".f16")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".s8")==0) && (strcmp(srcType, ".f16")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".s16")==0) && (strcmp(srcType, ".f16")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".s32")==0) && (strcmp(srcType, ".f16")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".s64")==0) && (strcmp(srcType, ".f16")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".u8")==0) && (strcmp(srcType, ".f32")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".u16")==0) && (strcmp(srcType, ".f32")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".u32")==0) && (strcmp(srcType, ".f32")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".u64")==0) && (strcmp(srcType, ".f32")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".s8")==0) && (strcmp(srcType, ".f32")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".s16")==0) && (strcmp(srcType, ".f32")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".s32")==0) && (strcmp(srcType, ".f32")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".s64")==0) && (strcmp(srcType, ".f32")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".u8")==0) && (strcmp(srcType, ".ff64")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".u16")==0) && (strcmp(srcType, ".ff64")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".u32")==0) && (strcmp(srcType, ".ff64")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".u64")==0) && (strcmp(srcType, ".ff64")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".s8")==0) && (strcmp(srcType, ".ff64")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".s16")==0) && (strcmp(srcType, ".ff64")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".s32")==0) && (strcmp(srcType, ".ff64")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".s64")==0) && (strcmp(srcType, ".ff64")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".f16")==0) && (strcmp(srcType, ".f16")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".f32")==0) && (strcmp(srcType, ".f32")==0)) - { output(modString); output("i"); } - if((strcmp(dstType, ".ff64")==0) && (strcmp(srcType, ".ff64")==0)) - { output(modString); output("i"); } - } - else if ( - (strcmp(modString, ".rni")==0 || strcmp(modString, ".rmi")==0 || - strcmp(modString, ".rpi")==0 || strcmp(modString, ".rzi")==0) - ) - { - output(modString); - } - else - { - output("\nUnknown mod:"); output(currentPiece->stringText); - assert(0); - } - - currentPiece = currentPiece->nextString; - } - } else { - printBaseModifiers(); - } - - // If one type modifier, duplicate to two; - int numModifiers = m_typeModifiers->getSize(); - stringListPiece* currentPiece = m_typeModifiers->getListStart(); - if( numModifiers == 1 ) { - output(currentPiece->stringText); - output(currentPiece->stringText); - } else { - printTypeModifiers(); - } - - printOperands(); - - output(";"); - - } - - } - else if(strcmp(m_base, "shl")==0) - { - printDefaultPtx(); - } - else if(strcmp(m_base, "add")==0) - { - printLabel(); - printPredicate(); - - output(m_base); - - printBaseModifiers(); - - currentPiece = m_typeModifiers->getListStart(); - for(int i=0; (i<m_typeModifiers->getSize())&&(currentPiece!=NULL); i++) - { - if(strcmp(currentPiece->stringText, ".b8")==0) - { - output(".u8"); - } - else if(strcmp(currentPiece->stringText, ".b16")==0) - { - output(".u16"); - } - else if(strcmp(currentPiece->stringText, ".b32")==0) - { - output(".u32"); - } - else if(strcmp(currentPiece->stringText, ".b64")==0) - { - output(".u64"); - } - else - { - output(currentPiece->stringText); - } - - currentPiece = currentPiece->nextString; - } - - printOperands(); - - output(";"); - } - else if(strcmp(m_base, "movsh")==0) - { - printLabel(); - printPredicate(); - - output("shl"); - - printBaseModifiers(); - - printTypeModifiers(); - - printOperands(); - - output(";"); - } - else if(strcmp(m_base, "mov")==0) - { - printLabel(); - printPredicate(); - - output("mov"); - - printBaseModifiers(); - - // If mov has two type modifiers, pick the smaller one for proper - // bit truncation - if( m_typeModifiers->getSize() == 2 ) { - std::string type1, type2, type; - int type1Size, type2Size; - stringListPiece* currentPiece = m_typeModifiers->getListStart(); - type1 = currentPiece->stringText; - type2 = currentPiece->nextString->stringText; - - type1Size = atoi(type1.substr(2, type1.size()-2).c_str()); - type2Size = atoi(type2.substr(2, type2.size()-2).c_str()); - type = (type1Size < type2Size) ? type1 : type2; - output(type.c_str()); - } else if( m_typeModifiers->getSize() == 1 ) { - printTypeModifiers(); - } else { - output("Error: unsupported number of type modifiers. "); - } - - printOperands(); - - output(";"); - } - else if(strcmp(m_base, "bar.sync")==0) - { - printLabel(); - printPredicate(); - - output(m_base); - - printBaseModifiers(); - - stringListPiece* currentPiece = m_typeModifiers->getListStart(); - currentPiece = m_typeModifiers->getListStart(); - for(int i=0; (i<m_typeModifiers->getSize())&&(currentPiece!=NULL); i++) - { - //output(currentPiece->stringText); - - currentPiece = currentPiece->nextString; - } - - printOperands(); - - output(";"); - m_opPerCycle = -1; - } - else if(strcmp(m_base, "mul")==0) - { - printDefaultPtx(); - - // opPerCycle - lower if a 32-bit integer mul - const char* dstType = m_typeModifiers->getListStart()->stringText; - if( strcmp(dstType, ".s32")==0 || strcmp(dstType, ".u32")==0 || strcmp(dstType, ".b32")==0 ) - m_opPerCycle = 2; - } - else if(strcmp(m_base, "mad24")==0) - { - printLabel(); - printPredicate(); - - output("mad24"); - - printBaseModifiers(); - - // Only output the destination operand type (first type modifier only) - stringListPiece* currentPiece = m_typeModifiers->getListStart(); - output(currentPiece->stringText); - - printOperands(); - // If 3 operands, this is a mac24 instruction, so add destination operand as 4th operand - int numOperands = m_operands->getSize(); - if(numOperands == 3) { - stringListPiece* currentPiece = m_operands->getListStart(); - output(", "); output(currentPiece->stringText); - } - - output(";"); - } - else if(strcmp(m_base, "mad24c1")==0) - { - output("nop; //"); - - printLabel(); - printPredicate(); - - output("mad24c1"); - - printBaseModifiers(); - - // Only output the destination operand type (first type modifier only) - stringListPiece* currentPiece = m_typeModifiers->getListStart(); - output(currentPiece->stringText); - - printOperands(); - // If 3 operands, this is a mac24 instruction, so add destination operand as 4th operand - int numOperands = m_operands->getSize(); - if(numOperands == 3) { - stringListPiece* currentPiece = m_operands->getListStart(); - output(", "); output(currentPiece->stringText); - } - - output(";"); - } - else if(strcmp(m_base, "set")==0) - { - printLabel(); - printPredicate(); - - output("set"); - - printBaseModifiers(); - - // If one type modifier, duplicate to two; - // if three type modifiers, remove third - int numModifiers = m_typeModifiers->getSize(); - stringListPiece* currentPiece = m_typeModifiers->getListStart(); - if( numModifiers == 1 ) { - output(currentPiece->stringText); - output(currentPiece->stringText); - } else if( numModifiers == 3) { - for(int i=0; (i<numModifiers-1)&&(currentPiece!=NULL); i++) - { - output(currentPiece->stringText); - currentPiece = currentPiece->nextString; - } - } - - printOperands(); - - output(";"); - } - else if(strcmp(m_base, "mad")==0) - { - printDefaultPtx(); - - // opPerCycle - lower if a 32-bit integer mad - const char* dstType = m_typeModifiers->getListStart()->stringText; - if( strcmp(dstType, ".s32")==0 || strcmp(dstType, ".u32")==0 || strcmp(dstType, ".b32")==0 ) - m_opPerCycle = 2; - } - else if(strcmp(m_base, "mul24")==0) - { - printLabel(); - printPredicate(); - - output("mul24"); - - printBaseModifiers(); - - // Only output the destination operand type (first type modifier only) - stringListPiece* currentPiece = m_typeModifiers->getListStart(); - output(currentPiece->stringText); - - printOperands(); - - output(";"); - } - else if(strcmp(m_base, "set?68?")==0) - { - // This actually takes absolute value of first source operand and is set.gt - printLabel(); - printPredicate(); - - output("set.gt.abs"); - - printBaseModifiers(); - - // If one type modifier, duplicate to two; - // if three type modifiers, remove third - int numModifiers = m_typeModifiers->getSize(); - stringListPiece* currentPiece = m_typeModifiers->getListStart(); - if( numModifiers == 1 ) { - output(currentPiece->stringText); - output(currentPiece->stringText); - } else if( numModifiers == 3) { - for(int i=0; (i<numModifiers-1)&&(currentPiece!=NULL); i++) - { - output(currentPiece->stringText); - currentPiece = currentPiece->nextString; - } - } - - printOperands(); - - output(";"); - - output(" //set?68?"); - } - else if(strcmp(m_base, "set?65?")==0) - { - // This actually takes absolute value of first source operand and is set.lt - printLabel(); - printPredicate(); - - output("set.lt.abs"); - - printBaseModifiers(); - - // If one type modifier, duplicate to two; - // if three type modifiers, remove third - int numModifiers = m_typeModifiers->getSize(); - stringListPiece* currentPiece = m_typeModifiers->getListStart(); - if( numModifiers == 1 ) { - output(currentPiece->stringText); - output(currentPiece->stringText); - } else if( numModifiers == 3) { - for(int i=0; (i<numModifiers-1)&&(currentPiece!=NULL); i++) - { - output(currentPiece->stringText); - currentPiece = currentPiece->nextString; - } - } - - printOperands(); - - output(";"); - - output(" //set?65?"); - } - else if(strcmp(m_base, "set?67?")==0) - { - // Change to set.gt - printLabel(); - printPredicate(); - - output("set.gt.abs"); - - printBaseModifiers(); - - // If one type modifier, duplicate to two; - // if three type modifiers, remove third - int numModifiers = m_typeModifiers->getSize(); - stringListPiece* currentPiece = m_typeModifiers->getListStart(); - if( numModifiers == 1 ) { - output(currentPiece->stringText); - output(currentPiece->stringText); - } else if( numModifiers == 3) { - for(int i=0; (i<numModifiers-1)&&(currentPiece!=NULL); i++) - { - output(currentPiece->stringText); - currentPiece = currentPiece->nextString; - } - } - - printOperands(); - - output(";"); - } - else if(strcmp(m_base, "set?13?")==0) - { - // Change to set.gt - printLabel(); - printPredicate(); - - output("set.neu"); - - printBaseModifiers(); - - // If one type modifier, duplicate to two; - // if three type modifiers, remove third - int numModifiers = m_typeModifiers->getSize(); - stringListPiece* currentPiece = m_typeModifiers->getListStart(); - if( numModifiers == 1 ) { - output(currentPiece->stringText); - output(currentPiece->stringText); - } else if( numModifiers == 3) { - for(int i=0; (i<numModifiers-1)&&(currentPiece!=NULL); i++) - { - output(currentPiece->stringText); - currentPiece = currentPiece->nextString; - } - } - - printOperands(); - - output(";"); - } - else if(strcmp(m_base, "rcp")==0) - { - printDefaultPtx(); - - m_opPerCycle = 2; - } - else if(strcmp(m_base, "pre.sin")==0) - { - printLabel(); - printPredicate(); - - output("nop;"); - output(" //"); - output("pre.sin"); - printBaseModifiers(); - printTypeModifiers(); - printOperands(); - - // m_opPerCycle - Ignore - m_opPerCycle = -1; - } - else if(strcmp(m_base, "sin")==0) - { - printDefaultPtx(); - - m_opPerCycle = 1; - } - else if(strcmp(m_base, "pre.ex2")==0) - { - printLabel(); - printPredicate(); - - output("ex2"); - printBaseModifiers(); - printTypeModifiers(); - printOperands(); - output(";"); - - output(" //"); - printDefaultPtx(); - - // m_opPerCycle - Ignore - m_opPerCycle = -1; - } - else if(strcmp(m_base, "ex2")==0) - { - printLabel(); - printPredicate(); - - output("nop;"); - output(" //"); - output("ex2"); - printBaseModifiers(); - printTypeModifiers(); - printOperands(); - - m_opPerCycle = 1; - } - else if(strcmp(m_base, "cos")==0) - { - printDefaultPtx(); - - m_opPerCycle = 1; - } - else if(strcmp(m_base, "lg2")==0) - { - printDefaultPtx(); - - m_opPerCycle = 2; - } - else if(strcmp(m_base, "rsqrt")==0) - { - printDefaultPtx(); - - m_opPerCycle = 2; - } - else if(strcmp(m_base, "mac")==0) - { - // Replace with mad by adding a 4th operand - printLabel(); - printPredicate(); - - output("mad"); - - printBaseModifiers(); - - printTypeModifiers(); - - // Print operands and then include destination (1st) operand as the 4th operand - printOperands(); - stringListPiece* currentPiece = m_operands->getListStart(); - output(","); output(" "); output(currentPiece->stringText); - - output(";"); - - // opPerCycle - lower if a 32-bit integer mac - const char* dstType = m_typeModifiers->getListStart()->stringText; - if( strcmp(dstType, ".s32")==0 || strcmp(dstType, ".u32")==0 || strcmp(dstType, ".b32")==0 ) - m_opPerCycle = 2; - } - else if(strcmp(m_base, "bra.label")==0) - { - printLabel(); - printPredicate(); - - output("bra"); - - printBaseModifiers(); - - printTypeModifiers(); - - printOperands(); - - output(";"); - } - else if(strcmp(m_base, "join.label")==0) - { - printLabel(); - printPredicate(); - - output("nop;"); - output(" //join.label"); - - // Ignore - m_opPerCycle = -1; - } - else if(strcmp(m_base, "nop.end")==0) - { - printLabel(); - printPredicate(); - - output("nop.exit;"); - - // Ignore - m_opPerCycle = -1; - } - else if(strcmp(m_base, "nop.join")==0) - { - printLabel(); - printPredicate(); - - output("nop;"); - - // Ignore - m_opPerCycle = -1; - } - else if(strcmp(m_base, "nop")==0) - { - //printLabel(); - //printPredicate(); - - output("nop;"); - } - else if(strcmp(m_base, "return")==0) - { - // ret instruction causes a deadlock bug in the simulator - // Temporary fix: branch to a dummy exit instruction that is added to the end of each - // entry with the label 'l_exit' - printLabel(); - printPredicate(); - - output("retp"); - //output("bra l_exit"); - - printBaseModifiers(); - - printTypeModifiers(); - - printOperands(); - - output(";"); - } - else if(strcmp(m_base, "and")==0) - { - printDefaultPtx(); - } - else if(strcmp(m_base, "andn")==0) - { - printDefaultPtx(); - } - else if(strcmp(m_base, "tex")==0) - { - printDefaultPtx(); - } - else if(strcmp(m_base, "xor")==0) - { - printDefaultPtx(); - } - else if(strcmp(m_base, "or")==0) - { - printDefaultPtx(); - } - else if(strcmp(m_base, "shr")==0) - { - printDefaultPtx(); - } - else if(strcmp(m_base, "subr")==0) - { - - // Replace with sub instruction with destination operands switched - printLabel(); - printPredicate(); - - output("sub"); - - printBaseModifiers(); - - printTypeModifiers(); - - // Switch source operands before printing - // Must be 3 operands, switch the last two; - if( m_operands->getSize() != 3) { - output("Error: subr instruction with number of operands other than 3.\n"); - assert(0); - } - const char* firstOperand = m_operands->getListStart()->stringText; - const char* secondOperand = m_operands->getListStart()->nextString->stringText; - const char* thirdOperand = m_operands->getListStart()->nextString->nextString->stringText; - output(" "); output(firstOperand); output(","); - output(" "); output(thirdOperand); output(","); - output(" "); output(secondOperand); - - - output(";"); - } - else if(strcmp(m_base, "sub")==0) - { - printDefaultPtx(); - } - else if(strcmp(m_base, "max")==0) - { - printDefaultPtx(); - } - else if(strcmp(m_base, "min")==0) - { - printDefaultPtx(); - } - else if(strcmp(m_base, "call.label")==0) - { - printLabel(); - printPredicate(); - - output("callp"); - - printBaseModifiers(); - - printTypeModifiers(); - - printOperands(); - - output(";"); - - // Ignore - m_opPerCycle = -1; - } - else if(strcmp(m_base, "not")==0) - { - printDefaultPtx(); - } - else if(strcmp(m_base, "delta")==0) - { - // This is a neg instruction - printLabel(); - printPredicate(); - - output("neg"); - - printBaseModifiers(); - - printTypeModifiers(); - - printOperands(); - - output(";"); - } - else if(strcmp(m_base, "break")==0) - { - printDefaultPtx(); - } - else if(strcmp(m_base, "breakaddr.label")==0) - { - printLabel(); - printPredicate(); - - output("breakaddr"); - - printBaseModifiers(); - - printTypeModifiers(); - - printOperands(); - - output(";"); - } - else if(strcmp(m_base, "inc")==0) - { - printDefaultPtx(); - } - else if(strcmp(m_base, "exch")==0) - { - printDefaultPtx(); - } - else if(strcmp(m_base, "cas")==0) - { - printDefaultPtx(); - } - else if(strcmp(m_base, "norn")==0) - { - printDefaultPtx(); - } - else if(strcmp(m_base, "addc")==0) - { - output("nop; //"); - printDefaultPtx(); - } - else if(strcmp(m_base, "orn")==0) - { - printDefaultPtx(); - } - else if(strcmp(m_base, "nandn")==0) - { - printDefaultPtx(); - } - else if(strcmp(m_base, "nxor")==0) - { - output("nop; //"); - printDefaultPtx(); - } - else if(strcmp(m_base, "sad")==0) - { - output("nop; //"); - printDefaultPtx(); - } - else if(strcmp(m_base, "op.13")==0) - { - output("nop; //"); - printDefaultPtx(); - } - else if(strcmp(m_base, "op.e5")==0) - { - output("nop; //"); - printDefaultPtx(); - } - else if(strcmp(m_base, "op.e6")==0) - { - output("nop; //"); - printDefaultPtx(); - } - else if(strcmp(m_base, "op.d0")==0) - { - output("nop; //"); - printDefaultPtx(); - } - - - else if(strcmp(m_base, "{")==0) - { - //output(m_base); - - // Ignore - m_opPerCycle = -1; - } - else if(strcmp(m_base, "}")==0) - { - //output(m_base); - - // Ignore - m_opPerCycle = -1; - } - else - { - output("Unknown Instruction: "); output(m_base); - assert(0); - } - -} |