From 46045592211f1999983755726263c60267f8ef2f Mon Sep 17 00:00:00 2001
From: Jimmy Kwa <jkwa@ece.ubc.ca>
Date: Thu, 6 Sep 2012 13:53:41 -0800
Subject: 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]
---
 cuobjdump_to_ptxplus/cuobjdumpInstList.cc | 38 +++++++++++++++++++++++++++++--
 1 file changed, 36 insertions(+), 2 deletions(-)

(limited to 'cuobjdump_to_ptxplus/cuobjdumpInstList.cc')

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() {
-- 
cgit v1.3