Fixing one of the problems discovered for bug 154 (internal).

  - Output of 16-bit integer multiplication (IMUL.U16) writes to the entire 32-bit register if a full register (i.e. not half register such as R1L) is specified as the destination.  Before, it was only storing the lower 16-bit of the product into the full register.
  - Modified function header to declare at least 4 predicates in cuobjdump_to_ptxplus.  This will mask out the problem with the predicate registers used before initialized.
  - Modified register value print out: now unsigned integers display their decimal values as well.

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

4 files changed, 27 insertions, 6 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");
 	}
diff --git a/src/cuda-sim/ptx_sim.cc b/src/cuda-sim/ptx_sim.cc
index 59479c8..05f0c6a 100644
--- a/src/cuda-sim/ptx_sim.cc
+++ b/src/cuda-sim/ptx_sim.cc
@@ -293,10 +293,10 @@ static void print_reg( FILE *fp, std::string name, ptx_reg_t value, symbol_table
    case S16_TYPE: fprintf(fp,".s16 %d\n", value.s16 ); break;
    case S32_TYPE: fprintf(fp,".s32 %d\n", value.s32 ); break;
    case S64_TYPE: fprintf(fp,".s64 %Ld\n", value.s64 ); break;
-   case U8_TYPE:  fprintf(fp,".u8  0x%02x\n", (unsigned) value.u8 );  break;
-   case U16_TYPE: fprintf(fp,".u16 0x%04x\n", (unsigned) value.u16 ); break;
-   case U32_TYPE: fprintf(fp,".u32 0x%08x\n", (unsigned) value.u32 ); break;
-   case U64_TYPE: fprintf(fp,".u64 0x%llx\n", value.u64 ); break;
+   case U8_TYPE:  fprintf(fp,".u8  %u [0x%02x]\n", value.u8, (unsigned) value.u8 );  break;
+   case U16_TYPE: fprintf(fp,".u16 %u [0x%04x]\n", value.u16, (unsigned) value.u16 ); break;
+   case U32_TYPE: fprintf(fp,".u32 %u [0x%08x]\n", value.u32, (unsigned) value.u32 ); break;
+   case U64_TYPE: fprintf(fp,".u64 %llu [0x%llx]\n", value.u64, value.u64 ); break;
    case F16_TYPE: fprintf(fp,".f16 %f [0x%04x]\n",  value.f16, (unsigned) value.u16 ); break;
    case F32_TYPE: fprintf(fp,".f32 %.15lf [0x%08x]\n",  value.f32, value.u32 ); break;
    case F64_TYPE: fprintf(fp,".f64 %.15le [0x%016llx]\n", value.f64, value.u64 ); break;