3 files changed, 23 insertions, 2 deletions

diff --git a/cuobjdump_to_ptxplus/cuobjdumpInst.cc b/cuobjdump_to_ptxplus/cuobjdumpInst.cc
index fac53b4..2bdac6d 100644
--- a/cuobjdump_to_ptxplus/cuobjdumpInst.cc
+++ b/cuobjdump_to_ptxplus/cuobjdumpInst.cc
@@ -721,6 +721,24 @@ void cuobjdumpInst::printCuobjdumpOperands()
 	}
 }
 
+void cuobjdumpInst::printCuobjdumpOutputModifiers(const char* defaultMod)
+{
+	std::list<std::string>::iterator typemod = m_typeModifiers->begin();
+	if (*typemod == ".U16" or *typemod == ".S16") {
+		std::list<std::string>::iterator dest_op = m_operands->begin(); 
+		std::string& destination = *dest_op; 
+		if (destination[destination.length()-1] == 'l') {
+			output(".lo");  // write to the lower 16-bits 
+		} else if (destination[destination.length()-1] == 'h') {
+			output(".hi");  // write to the upper 16-bits 
+		} else {
+			output(".wide");  // write to the whole 32-bits 
+		}
+		return; 
+	}
+	output(defaultMod);  // default output modifier for mul 
+}
+
 std::string int_default_mod () { return ".u32" ;}
 
 
@@ -1309,7 +1327,8 @@ void cuobjdumpInst::printCuobjdumpPtxPlus(std::list<std::string> labelList, std:
 	else if(m_base == "IMUL" || m_base == "IMUL32I")
 	{
 		printCuobjdumpPredicate();
-		output("mul.lo");
+		output("mul");
+      printCuobjdumpOutputModifiers(".lo"); 
 		printCuobjdumpBaseModifiers();
 
 		if(m_typeModifiers->size() == 0)
diff --git a/cuobjdump_to_ptxplus/cuobjdumpInst.h b/cuobjdump_to_ptxplus/cuobjdumpInst.h
index 3a09e64..c48d7a5 100644
--- a/cuobjdump_to_ptxplus/cuobjdumpInst.h
+++ b/cuobjdump_to_ptxplus/cuobjdumpInst.h
@@ -72,6 +72,7 @@ public:
 	void printCuobjdumpLabel(std::list<std::string> labelList);
 	void printCuobjdumpPredicate();
 	void printCuobjdumpTypeModifiers();
+	void printCuobjdumpOutputModifiers(const char* defaultMod);
 	void printCuobjdumpBaseModifiers();
 	void printCuobjdumpOperand(std::string currentPiece, std::string operandDelimiter, std::string base);
 	void printCuobjdumpOperandlohi(std::string op);
diff --git a/cuobjdump_to_ptxplus/cuobjdumpInstList.cc b/cuobjdump_to_ptxplus/cuobjdumpInstList.cc
index 2c6a4e8..fabc396 100644
--- a/cuobjdump_to_ptxplus/cuobjdumpInstList.cc
+++ b/cuobjdump_to_ptxplus/cuobjdumpInstList.cc
@@ -211,7 +211,8 @@ void cuobjdumpInstList::printPredNames(cuobjdumpEntry entry)
 {
 	if( entry.m_largestPredIndex >= 0) {
 		char out[30];
-		sprintf(out, "\t.reg .pred $p<%d>;", entry.m_largestPredIndex+1);
+		// there is at least 4 predicates for GT200, possibly more in Fermi 
+		sprintf(out, "\t.reg .pred $p<%d>;", std::max(entry.m_largestPredIndex+1, 4)); 
 		output(out);
 		output("\n");
 	}