Re-adds simple version of combining VPM setups

This version will only combine writing of same setup values, where possible. The full version is also removed, since it will anyway become obsolete with VPM cached memory (see #113). Effects (test-emulator): Instructions: 52511 to 49793 (-5%) Cycles: 644891 to 641680 (-0.5%) Total time (in ms): 62869 to 58456 (-7%)
doe300 · Dec 21, 2018 · 7bad8c0 · 7bad8c0
1 parent c873760
commit 7bad8c0
Show file tree

Hide file tree

Showing 7 changed files with 133 additions and 557 deletions.
diff --git a/src/normalization/AddressCalculation.cpp b/src/normalization/AddressCalculation.cpp
@@ -16,115 +16,6 @@ using namespace vc4c::intermediate;
 using namespace vc4c::normalization;
 using namespace vc4c::operators;
 
-static BaseAndOffset findOffset(const Value& val)
-{
-    if(!val.hasLocal())
-        return BaseAndOffset();
-    const LocalUser* writer = val.getSingleWriter();
-    if(writer != nullptr)
-    {
-        const Optional<Value> offset = writer->precalculate(8);
-        if(offset && offset->isLiteralValue())
-        {
-            return BaseAndOffset(NO_VALUE, offset->getLiteralValue()->signedInt());
-        }
-    }
-    return BaseAndOffset();
-}
-
-BaseAndOffset normalization::findBaseAndOffset(const Value& val)
-{
-    // TODO add support for offsets via getlocal/global_id, etc.
-    // need to the set base to addr + offset and the offset to the offset of the offset (e.g. param[get_local_id(0) +
-    // 7])  but how to determine?
-    if(!val.hasLocal())
-        return BaseAndOffset();
-    if(val.local()->is<Parameter>() || val.local()->is<Global>() || val.local()->is<StackAllocation>())
-        return BaseAndOffset(val, 0);
-
-    // follow the references
-    const Local* ref = val.local()->getBase(false);
-    if(ref != val.local())
-        return findBaseAndOffset(ref->createReference());
-    if(val.local()->reference.first != nullptr && val.local()->reference.second != ANY_ELEMENT)
-        return BaseAndOffset(val.local()->reference.first->createReference(), val.local()->reference.second);
-
-    const auto writers = val.local()->getUsers(LocalUse::Type::WRITER);
-    if(writers.size() != 1)
-        return BaseAndOffset();
-
-    // The reader can be one of several valid cases:
-    // 1. a move from another local -> need to follow the move
-    if(dynamic_cast<const MoveOperation*>(*writers.begin()) != nullptr)
-        return findBaseAndOffset(dynamic_cast<const MoveOperation*>(*writers.begin())->getSource());
-    const auto& args = (*writers.begin())->getArguments();
-    // 2. an addition with a local and a literal -> the local is the base, the literal the offset
-    if(dynamic_cast<const Operation*>((*writers.begin())) != nullptr &&
-        dynamic_cast<const Operation*>((*writers.begin()))->op == OP_ADD && args.size() == 2 &&
-        std::any_of(args.begin(), args.end(), [](const Value& arg) -> bool { return arg.hasLocal(); }) &&
-        std::any_of(
-            args.begin(), args.end(), [](const Value& arg) -> bool { return arg.getLiteralValue().has_value(); }))
-    {
-        return BaseAndOffset(
-            std::find_if(args.begin(), args.end(), [](const Value& arg) -> bool { return arg.hasLocal(); })
-                ->local()
-                ->getBase(false)
-                ->createReference(),
-            static_cast<int32_t>((*std::find_if(args.begin(), args.end(),
-                                      [](const Value& arg) -> bool { return arg.getLiteralValue().has_value(); }))
-                                     .getLiteralValue()
-                                     ->signedInt() /
-                val.type.getElementType().getPhysicalWidth()));
-    }
-
-    // 3. an addition with two locals -> one is the base, the other the calculation of the literal
-    if(dynamic_cast<const Operation*>((*writers.begin())) != nullptr &&
-        dynamic_cast<const Operation*>((*writers.begin()))->op == OP_ADD && args.size() == 2 &&
-        std::all_of(args.begin(), args.end(), [](const Value& arg) -> bool { return arg.hasLocal(); }))
-    {
-        const auto offset0 = findOffset(args[0]);
-        const auto offset1 = findOffset(args[1]);
-        if(offset0.offset && args[1].hasLocal())
-            return BaseAndOffset(args[1].local()->getBase(false)->createReference(),
-                static_cast<int32_t>(offset0.offset.value() / val.type.getElementType().getPhysicalWidth()));
-        if(offset1.offset && args[0].hasLocal())
-            return BaseAndOffset(args[0].local()->getBase(false)->createReference(),
-                static_cast<int32_t>(offset1.offset.value() / val.type.getElementType().getPhysicalWidth()));
-    }
-    /*
-        if(writers.size() == 1)
-        {
-            // couldn't find literal offset for any direct base, try with arbitrary values
-            ref = val.local->getBase(true);
-            Optional<Value> offset = NO_VALUE;
-            for(const auto& arg : (*writers.begin())->getArguments())
-            {
-                if(ref != nullptr && arg.hasLocal() && arg.local->getBase(false) == ref)
-                    // skip finding the same base again
-                    continue;
-                auto tmp = findBaseAndOffset(arg);
-                if(tmp.base && tmp.base->local->getBase(true) == ref && tmp.offset.is(0))
-                    // this parameter is the base itself, is already handled
-                    continue;
-                // TODO how to combine the offsets?
-                // TODO also need to handle non-addition of offsets (e.g. ptr = base + (offset + size * i))
-                logging::debug() << "Found offset of " << tmp.base.to_string() << " + "
-                                 << (tmp.offset ? tmp.offset.value() : -1) << logging::endl;
-                logging::debug() << "Found offset of " << tmp.base.to_string() << " with expression: "
-                                 << vc4c::Expression::createExpression(*(*writers.begin())).to_string() <<
-       logging::endl;
-            }
-            // TODO why is this called twice? The whole function, from outside
-            logging::debug() << "Found base and non-literal offset: " << ref->to_string() << " - " << offset.to_string()
-                             << logging::endl;
-            if(ref && (ref->residesInMemory() || (ref->is<Parameter>() && ref->type.isPointerType())))
-                return BaseAndOffset(ref->createReference(), {});
-        }
-    */
-
-    return BaseAndOffset();
-}
-
 MemoryType normalization::toMemoryType(periphery::VPMUsage usage)
 {
     switch(usage)

diff --git a/src/normalization/AddressCalculation.h b/src/normalization/AddressCalculation.h
@@ -16,21 +16,6 @@ namespace vc4c
 {
     namespace normalization
     {
-        struct BaseAndOffset
-        {
-            Optional<Value> base;
-            Optional<int32_t> offset;
-
-            explicit BaseAndOffset() : base(NO_VALUE), offset{} {}
-
-            BaseAndOffset(const Optional<Value>& base, Optional<int32_t> offset) : base(base), offset(offset) {}
-        };
-
-        /*
-         * Returns the base memory address for the given value as well as the offset the value has to this location
-         */
-        BaseAndOffset findBaseAndOffset(const Value& val);
-
         enum class MemoryType
         {
             // lower the value into a register and replace all loads with moves