Several changes to ptxplus towards getting the sphyraena benchmark to work. Add .false modifier to cuobjdump_to_ptxplus. It isn't supported in gpgpusim yet since we don't know what it does. Fixed most of the issues with global (const14) variable not being declared properly. Added "BRX" instruction to cuobjdump_to_ptxplus and some support for it in gpgpusim. There are other issues with it that still need to be worked out. Added support for translating IADD.CARRY sass instruction into ptxplus addp instructions. Fixed a bug with generating ".half" modifiers on instructions so the ptxplus instruction addresses should match sass addresses now. Separated ssy instructions from nop. However, they are still simulated as nop in gpgpusim. Sometime cuobjdump generate a blank link that takes up an address slot. This was changed to a nop.

[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 14001]

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	{ }