summaryrefslogtreecommitdiff
path: root/cuobjdump_to_ptxplus
diff options
context:
space:
mode:
Diffstat (limited to 'cuobjdump_to_ptxplus')
-rw-r--r--cuobjdump_to_ptxplus/cuobjdumpInst.cc75
-rw-r--r--cuobjdump_to_ptxplus/cuobjdumpInstList.cc38
-rw-r--r--cuobjdump_to_ptxplus/cuobjdumpInstList.h5
-rw-r--r--cuobjdump_to_ptxplus/elf.l21
-rw-r--r--cuobjdump_to_ptxplus/elf.y23
-rw-r--r--cuobjdump_to_ptxplus/sass.l5
-rw-r--r--cuobjdump_to_ptxplus/sass.y17
7 files changed, 160 insertions, 24 deletions
diff --git a/cuobjdump_to_ptxplus/cuobjdumpInst.cc b/cuobjdump_to_ptxplus/cuobjdumpInst.cc
index a18bd04..c74d2d8 100644
--- a/cuobjdump_to_ptxplus/cuobjdumpInst.cc
+++ b/cuobjdump_to_ptxplus/cuobjdumpInst.cc
@@ -280,6 +280,8 @@ void cuobjdumpInst::printCuobjdumpPredicate()
output(".sf");
} else if(modString3 == ".carry") {
output(".cf");
+ } else if(modString3 == ".false") {
+ output(".false"); //TODO: Need to find out what this is.
} else {
output(modString3);
}
@@ -496,7 +498,7 @@ void cuobjdumpInst::printCuobjdumpOperand(std::string currentPiece, std::string
}
// Make it lower case
- if(mod.substr(0,9)!= "constant1")
+ if(mod.substr(0,9)!= "constant1" && mod.substr(0,9) != "varglobal")
for(unsigned i=0; i<mod.length(); i++)
{
mod[i] = tolower(mod[i]);
@@ -712,7 +714,11 @@ void cuobjdumpInst::printCuobjdumpOperand(std::string currentPiece, std::string
output(outputhex.c_str());
} else if(mod.substr(0,3) == "l0x") { //label
output(mod.c_str());
+ } else if(mod.substr(0,9) == "varglobal") { //global variable
+ output(mod.substr(9).c_str());
} else {//variable name
+ printf("Unrecognized Operand: %s\n", mod.c_str());
+ assert(0);
output(mod.c_str());
}
}
@@ -794,6 +800,20 @@ void cuobjdumpInst::printCuobjdumpPtxPlus(std::list<std::string> labelList, std:
printCuobjdumpOperands();
output(";");
}
+ else if(m_base == "BRX")
+ {
+ printCuobjdumpPredicate();
+ output("brx");
+ printCuobjdumpBaseModifiers();
+
+ if(m_typeModifiers->size() == 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(m_base == "CAL")
{
printCuobjdumpPredicate();
@@ -1180,7 +1200,7 @@ void cuobjdumpInst::printCuobjdumpPtxPlus(std::list<std::string> labelList, std:
printCuobjdumpOperands();
output(";");
}
- else if(m_base == "IADD.CARRY"){
+ else if( m_base.find("IADD.CARRY") == 0){ //searches for IADD.CARRY at the start to match IADD.CARRY{numeric}
std::string pred = "C0";
pred[1] = m_base[10];
this->addOperand(pred.c_str());
@@ -1196,7 +1216,7 @@ void cuobjdumpInst::printCuobjdumpPtxPlus(std::list<std::string> labelList, std:
printCuobjdumpOperands();
output(";");
}
- else if(m_base == "IADD")
+ else if(m_base == "IADD" || m_base == "IADD32I")
{
printCuobjdumpPredicate();
output("add");
@@ -1210,7 +1230,7 @@ void cuobjdumpInst::printCuobjdumpPtxPlus(std::list<std::string> labelList, std:
printCuobjdumpOperands();
output(";");
}
- else if(m_base == "IADD32" || m_base == "IADD32I")
+ else if(m_base == "IADD32")
{
printCuobjdumpPredicate();
output("add.half");
@@ -1347,6 +1367,42 @@ void cuobjdumpInst::printCuobjdumpPtxPlus(std::list<std::string> labelList, std:
printCuobjdumpOperands();
output(";");
}
+ else if(m_base == "IMUL.HI.U24.U24")
+ {
+ printCuobjdumpPredicate();
+ output("mul24.hi");
+ printCuobjdumpBaseModifiers();
+
+ if(m_typeModifiers->size() == 0)
+ {
+ output(".u32"); //TODO: setting default type modifier but I'm not sure if this is right.
+ }
+ else if(m_typeModifiers->size() == 2)
+ {
+ std::string type1, type2, type;
+ int type1Size, type2Size;
+ char tempString[5];
+
+ std::list<std::string>::iterator curr = m_typeModifiers->begin();
+ type1 = (*curr).c_str();
+ curr++;
+ type2 = (*curr).c_str();
+
+ 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(tempString[1] >= 'A' && tempString[1] <= 'Z')
+ tempString[1] += 32;
+ output(tempString);
+ }
+ else
+ {
+ printCuobjdumpTypeModifiers();
+ }
+ printCuobjdumpOperands();
+ output(";");
+ }
else if(m_base == "IMUL" || m_base == "IMUL32I")
{
printCuobjdumpPredicate();
@@ -1550,7 +1606,7 @@ void cuobjdumpInst::printCuobjdumpPtxPlus(std::list<std::string> labelList, std:
printCuobjdumpOperands();
output(";");
}
- else if(m_base == "NOP" || m_base == "SSY")
+ else if(m_base == "NOP")
{
printCuobjdumpPredicate();
output("nop");
@@ -1831,6 +1887,15 @@ void cuobjdumpInst::printCuobjdumpPtxPlus(std::list<std::string> labelList, std:
printCuobjdumpOperands();
output(";");
}
+ else if( m_base == "SSY")
+ {
+ printCuobjdumpPredicate();
+ output("ssy");
+ printCuobjdumpBaseModifiers();
+ printCuobjdumpTypeModifiers();
+ printCuobjdumpOperands();
+ output(";");
+ }
else if(m_base == "LD")
{
printCuobjdumpPredicate();
diff --git a/cuobjdump_to_ptxplus/cuobjdumpInstList.cc b/cuobjdump_to_ptxplus/cuobjdumpInstList.cc
index d3ebd9a..32834c7 100644
--- a/cuobjdump_to_ptxplus/cuobjdumpInstList.cc
+++ b/cuobjdump_to_ptxplus/cuobjdumpInstList.cc
@@ -225,7 +225,7 @@ void cuobjdumpInstList::printOutOfBoundRegisters(cuobjdumpEntry entry)
if( entry.m_reg124 == true ) {
output("\n");
output("\t.reg .u32 $r124;\n");
- output("\tmov.u32 $r124, 0x00000000;\n");
+ // output("\tmov.u32 $r124, 0x00000000;\n");
}
if( entry.m_oreg127 == true) {
output("\n");
@@ -392,7 +392,7 @@ void cuobjdumpInstList::addCuobjdumpMemoryOperand(std::string mem, int memType)
std::list<globalMemory>::iterator g;
for(g=m_globalMemoryList.begin(); g!=m_globalMemoryList.end(); ++g) {
if(g->offset == offset) {
- mem = g->name;
+ mem = "varglobal" + g->name;
found = true;
break;
}
@@ -584,6 +584,40 @@ void cuobjdumpInstList::setLocalMemoryMap(const char* kernelname, int index){
kernellmemmap[kernel] = index;
}
+void cuobjdumpInstList::setglobalVarShndx(const char* shndx){
+ m_globalVarShndx = atoi(shndx);
+}
+
+int cuobjdumpInstList::getglobalVarShndx(){
+ return m_globalVarShndx;
+}
+
+void cuobjdumpInstList::addGlobalMemoryID(const char* bytes, const char* name){
+ globalMemory globalMemID;
+ //globalMemID.offset = atoi(index)/4;
+ globalMemID.bytes = atoi(bytes);
+ globalMemID.name = name;
+
+ m_globalMemoryList.push_back(globalMemID);
+}
+
+void cuobjdumpInstList::updateGlobalMemoryID(const char* offset, const char* name){
+ bool found = false;
+ std::list<globalMemory>::iterator g;
+ for(g=m_globalMemoryList.begin(); g!=m_globalMemoryList.end(); ++g) {
+ if(g->name.compare(name) == 0) {
+ g->offset = atoi(offset)/4;
+ found = true;
+ break;
+ }
+ }
+
+ if(!found) {
+ printf("Could not find a global memory with this offset in: %s\n", name);
+ output("Could not find a global memory with this offset.\n");
+ assert(0);
+ }
+}
//NOT USED
void cuobjdumpInstList::reverseConstMemory() {
diff --git a/cuobjdump_to_ptxplus/cuobjdumpInstList.h b/cuobjdump_to_ptxplus/cuobjdumpInstList.h
index d07e72d..da8f8b0 100644
--- a/cuobjdump_to_ptxplus/cuobjdumpInstList.h
+++ b/cuobjdump_to_ptxplus/cuobjdumpInstList.h
@@ -116,6 +116,7 @@ protected:
std::list<localMemory> m_localMemoryList;
std::list<std::string> m_realTexList; // Stores the real names of tex variables
std::list<constMemoryPtr> m_constMemoryPtrList;
+ int m_globalVarShndx; //records shndx value of global variables in the elf file SYMTAB
// Functions:
std::string parseCuobjdumpPredicate(std::string pred);
@@ -158,6 +159,10 @@ public:
void addConstMemoryPtr(const char* bytes, const char* offset, const char* name);
void setConstMemoryMap(const char* kernelname, int index);
void setLocalMemoryMap(const char* kernelname, int index);
+ void setglobalVarShndx(const char* shndx);
+ int getglobalVarShndx();
+ void addGlobalMemoryID(const char* bytes, const char* name);
+ void updateGlobalMemoryID(const char* offset, const char* name);
void reverseConstMemory();
void addEntryLocalMemory(int value, int entryIndex);
void readOtherConstMemoryFromBinFile(std::string binString); // read in constant memory from bin file
diff --git a/cuobjdump_to_ptxplus/elf.l b/cuobjdump_to_ptxplus/elf.l
index 61c4d02..956550c 100644
--- a/cuobjdump_to_ptxplus/elf.l
+++ b/cuobjdump_to_ptxplus/elf.l
@@ -43,6 +43,7 @@ void elf_error(const char*);
%x cmem
%x symtab
+%x cmem14
whitespace [ \t]
newline [\n]
@@ -52,23 +53,27 @@ ident [[:alnum:]_\.]
hexdigit [[:digit:]abcdef]
%%
-<cmem,INITIAL>\.nv\.constant1\.{notnewline}*{newline} {
+<cmem,cmem14,INITIAL>\.nv\.constant1\.{notnewline}*{newline} {
BEGIN(cmem);
yylval.string_value = strdup(yytext);
return C1BEGIN;
}
-<cmem,INITIAL>\.nv\.local{notnewline}*{newline} {
+<cmem,cmem14,INITIAL>\.nv\.local{notnewline}*{newline} {
BEGIN(INITIAL);
yylval.string_value = strdup(yytext);
return LOCALMEM;
}
-<cmem,INITIAL>\.nv\.constant0{notnewline}*{newline} {
+<cmem,cmem14,INITIAL>\.nv\.constant0{notnewline}*{newline} {
BEGIN(cmem);
return C0BEGIN;
}
+<cmem,cmem14,INITIAL>\.section" "\.rel\.nv\.constant14{notnewline}*{newline} {
+ BEGIN(cmem14);
+ return C14BEGIN;
+}
\.section" "\.symtab\tSYMTAB{newline} {
BEGIN(symtab);
@@ -79,7 +84,7 @@ hexdigit [[:digit:]abcdef]
//These are the lines before the symtab
}
-<cmem,INITIAL>\.{notnewline}*{newline} {
+<cmem,cmem14,INITIAL>\.{notnewline}*{newline} {
BEGIN(INITIAL);
}
@@ -102,16 +107,16 @@ hexdigit [[:digit:]abcdef]
return STHEADER;
}
-<symtab>{whitespace}+ {}
+<symtab,cmem14>{whitespace}+ {}
-<symtab>{newline} {}
+<symtab,cmem14>{newline} {}
-<symtab>{hexdigit}+ {
+<symtab,cmem14>{hexdigit}+ {
yylval.string_value = strdup(yytext);
return NUMBER;
}
-<symtab>{ident}+ {
+<symtab,cmem14>{ident}+ {
yylval.string_value = strdup(yytext);
return IDENTIFIER;
}
diff --git a/cuobjdump_to_ptxplus/elf.y b/cuobjdump_to_ptxplus/elf.y
index f521518..fe471da 100644
--- a/cuobjdump_to_ptxplus/elf.y
+++ b/cuobjdump_to_ptxplus/elf.y
@@ -43,7 +43,7 @@ bool lastcmem = false;// false = constrant0, true = constant1
%union {
char* string_value;
}
-%token <string_value> C1BEGIN CMEMVAL SPACE2 C0BEGIN STBEGIN STHEADER
+%token <string_value> C1BEGIN C14BEGIN CMEMVAL SPACE2 C0BEGIN STBEGIN STHEADER
%token <string_value> NUMBER HEXNUMBER IDENTIFIER LOCALMEM
%%
@@ -58,8 +58,15 @@ stcontent : stcontent stline
| stline;
stline : NUMBER NUMBER NUMBER NUMBER NUMBER NUMBER IDENTIFIER {
+ if (strcmp($7, ".nv.global")==0) {
+ g_instList->setglobalVarShndx($6);
+ }
if (strcmp($4, "11")==0) {
- g_instList->addConstMemoryPtr($2, $3, $7);
+ if(atoi($6) == g_instList->getglobalVarShndx()) {
+ g_instList->addGlobalMemoryID($3, $7);
+ } else {
+ g_instList->addConstMemoryPtr($2, $3, $7);
+ }
}
}
| NUMBER NUMBER NUMBER NUMBER NUMBER NUMBER {}
@@ -67,6 +74,7 @@ stline : NUMBER NUMBER NUMBER NUMBER NUMBER NUMBER IDENTIFIER {
program : program cmemsection
| program localmemsection
+ | program cmem14section
| {
g_instList->setKernelCount(cmemcount-1);
};
@@ -96,4 +104,15 @@ cmemvals : cmemvals CMEMVAL SPACE2 {
g_instList->addConstMemoryValue($3);
}
| ;
+
+cmem14section : C14BEGIN c14content
+ | C14BEGIN
+ ;
+
+c14content : c14content c14line
+ | c14line;
+
+c14line : NUMBER IDENTIFIER IDENTIFIER {
+ g_instList->updateGlobalMemoryID($1, $2);
+ };
%%
diff --git a/cuobjdump_to_ptxplus/sass.l b/cuobjdump_to_ptxplus/sass.l
index b5b0532..22d9e2f 100644
--- a/cuobjdump_to_ptxplus/sass.l
+++ b/cuobjdump_to_ptxplus/sass.l
@@ -80,6 +80,7 @@ fltliteral [0-9]+\.[0-9]+
"BAR.ARV.WAIT b0, 0xfff" yylval.string_value = strdup(yytext); return BAR;
ADA yylval.string_value = strdup(yytext); return ADA;
BRA yylval.string_value = strdup(yytext); return BRA;
+BRX yylval.string_value = strdup(yytext); return BRX;
CAL yylval.string_value = strdup(yytext); return CAL;
COS yylval.string_value = strdup(yytext); return COS;
DADD yylval.string_value = strdup(yytext); return DADD;
@@ -122,11 +123,12 @@ IMAD32I yylval.string_value = strdup(yytext); return IMAD32I;
IMUL yylval.string_value = strdup(yytext); return IMUL;
IMUL\.U24\.U24 yylval.string_value = strdup(yytext); return IMUL24;
IMUL\.S24\.S24 yylval.string_value = strdup(yytext); return IMULS24;
+IMUL\.HI\.U24\.U24 yylval.string_value = strdup(yytext); return IMUL24H;
IMUL32 yylval.string_value = strdup(yytext); return IMUL32;
IMUL32\.S24\.S24 yylval.string_value = strdup(yytext); return IMUL32S24;
IMUL32\.U24\.U24 yylval.string_value = strdup(yytext); return IMUL32U24;
-IMUL32I\.S24\.S24 yylval.string_value = strdup(yytext); return IMUL32IS24;
IMUL32I yylval.string_value = strdup(yytext); return IMUL32I;
+IMUL32I\.S24\.S24 yylval.string_value = strdup(yytext); return IMUL32IS24;
IMUL32I\.U24\.U24 yylval.string_value = strdup(yytext); return IMUL32I24;
ISET yylval.string_value = strdup(yytext); return ISET;
LG2 yylval.string_value = strdup(yytext); return LG2;
@@ -207,6 +209,7 @@ NEU yylval.string_value = strdup(yytext); return NEU;
/* comparison operators */
\.EQ yylval.string_value = strdup(yytext); return DOTEQ;
\.EQU yylval.string_value = strdup(yytext); return DOTEQU;
+\.FALSE yylval.string_value = strdup(yytext); return DOTFALSE;
\.GE yylval.string_value = strdup(yytext); return DOTGE;
\.GEU yylval.string_value = strdup(yytext); return DOTGEU;
\.GT yylval.string_value = strdup(yytext); return DOTGT;
diff --git a/cuobjdump_to_ptxplus/sass.y b/cuobjdump_to_ptxplus/sass.y
index 3706ad8..1cb6388 100644
--- a/cuobjdump_to_ptxplus/sass.y
+++ b/cuobjdump_to_ptxplus/sass.y
@@ -48,10 +48,10 @@ cuobjdumpInst *instEntry;
}
%token <string_value> BAR
-%token <string_value> ADA AND ANDS BRA CAL COS DADD DMIN DMAX DFMA DMUL EX2 F2F F2I FADD
+%token <string_value> ADA AND ANDS BRA BRX CAL COS DADD DMIN DMAX DFMA DMUL EX2 F2F F2I FADD
%token <string_value> FADD32 FADD32I FMAD FMAD32I FMUL FMUL32 FMUL32I FSET DSET G2R
%token <string_value> GLD GST I2F I2I IADD IADD32 IADD32I IMAD ISAD IMAD24 IMAD32I IMAD32 IADDCARRY
-%token <string_value> IMUL IMUL24 IMULS24 IMUL32 IMUL32S24 IMUL32U24 IMUL32I IMUL32I24 IMUL32IS24
+%token <string_value> IMUL IMUL24 IMUL24H IMULS24 IMUL32 IMUL32S24 IMUL32U24 IMUL32I IMUL32I24 IMUL32IS24
%token <string_value> ISET LG2 LLD LST MOV MOV32 MVC MVI NOP NOT NOTS OR ORS
%token <string_value> R2A R2G R2GU16U8 RCP RCP32 RET RRO RSQ SIN SHL SHR SSY XOR XORS
%token <string_value> S2R SASS_LD STS LDS SASS_ST IMIN IMAX A2R FMAX FMIN TEX TEX32 C2R EXIT
@@ -62,7 +62,7 @@ cuobjdumpInst *instEntry;
%token <string_value> DOTTRUNC DOTCEIL DOTFLOOR DOTIR DOTUN DOTNODEP DOTSAT DOTANY DOTALL
%token <string_value> DOTF16 DOTF32 DOTF64 DOTS8 DOTS16 DOTS32 DOTS64 DOTS128 DOTU8 DOTU16 DOTU32 DOTU24 DOTU64
%token <string_value> DOTHI DOTNOINC
-%token <string_value> DOTEQ DOTEQU DOTGE DOTGEU DOTGT DOTGTU DOTLE DOTLEU DOTLT DOTLTU DOTNE DOTNEU DOTNSF DOTSF DOTCARRY
+%token <string_value> DOTEQ DOTEQU DOTFALSE DOTGE DOTGEU DOTGT DOTGTU DOTLE DOTLEU DOTLT DOTLTU DOTNE DOTNEU DOTNSF DOTSF DOTCARRY
%token <string_value> DOTCC DOTX DOTE DOTRED DOTPOPC
%token <string_value> REGISTER REGISTERLO REGISTERHI OFFSETREGISTER
%token <string_value> PREDREGISTER PREDREGISTER2 PREDREGISTER3 SREGISTER
@@ -121,9 +121,13 @@ statementList : statementList statement NEWLINE { debug_print("\n"); }
| NEWLINE {}
;
-statement : { instEntry = new cuobjdumpInst(); } instructionLabel instructionHex assemblyInstruction
+statement : { instEntry = new cuobjdumpInst(); } instructionLabel statementend
;
+statementend : instructionHex assemblyInstruction
+ | /*blank*/ {instEntry->setBase("NOP"); g_instList->add(instEntry); debug_print("NOP");}
+ ;
+
instructionHex : INSTHEX
;
@@ -157,11 +161,11 @@ baseInstruction : simpleInstructions { debug_print($1); instEntry->setBase($1);
| pbkInstruction
;
-simpleInstructions : ADA | AND | ANDS | COS | DADD | DMIN | DMAX | DFMA | DMUL | EX2 | F2F
+simpleInstructions : ADA | AND | ANDS | BRX | COS | DADD | DMIN | DMAX | DFMA | DMUL | EX2 | F2F
| F2I | FADD | FADD32 | FADD32I | FMAD | FMAD32I | FMUL
| FMUL32 | FMUL32I | FSET | DSET | G2R | GLD | GST | I2F | I2I
| IADD | IADD32 | IADD32I | IMAD | ISAD | IMAD24 | IMAD32I | IMAD32 | IMUL
- | IMUL24 | IMULS24 | IMUL32 | IMUL32S24 | IMUL32I | IMUL32I24 | IMUL32IS24
+ | IMUL24 | IMUL24H | IMULS24 | IMUL32 | IMUL32S24 | IMUL32I | IMUL32I24 | IMUL32IS24
| IMUL32U24
| ISET | LG2 | LLD | LST | MOV | MOV32 | MVC | MVI | NOP
| NOT | NOTS | OR | ORS | R2A | R2G | R2GU16U8 | RCP | RCP32 | RET | RRO
@@ -407,6 +411,7 @@ preOperand : EX2 { debug_print($1); g_instList->getListEnd().addBaseModifier("ex
predicateModifier : DOTEQ { }
| DOTEQU { }
+ | DOTFALSE { }
| DOTGE { }
| DOTGEU { }
| DOTGT { }
generated by cgit v1.3 (git 2.53.0) at 2026-07-10 22:44:21 +0000