From bb96888d31d37760c86fa1aed3bbf203792185e7 Mon Sep 17 00:00:00 2001 From: Wilson Fung Date: Sat, 30 Jun 2012 11:15:07 -0800 Subject: 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] --- cuobjdump_to_ptxplus/cuobjdumpInst.cc | 21 ++++++++++++++++++++- cuobjdump_to_ptxplus/cuobjdumpInst.h | 1 + cuobjdump_to_ptxplus/cuobjdumpInstList.cc | 3 ++- 3 files changed, 23 insertions(+), 2 deletions(-) (limited to 'cuobjdump_to_ptxplus') 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::iterator typemod = m_typeModifiers->begin(); + if (*typemod == ".U16" or *typemod == ".S16") { + std::list::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 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 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"); } -- cgit v1.3