text:
// Copyright (C) 2005, 2006 The Trustees of Indiana University.

// Use, modification and distribution is subject to the Boost Software
// License, Version 1.0. (See accompanying file LICENSE_1_0.txt or copy at
// http://www.boost.org/LICENSE_1_0.txt)

//  Authors: Douglas Gregor
//           Andrew Lumsdaine
#define PBGL_ACCOUNTING

#include <boost/graph/use_mpi.hpp>
#include <boost/config.hpp>
#include <boost/throw_exception.hpp>
#include <boost/graph/distributed/boman_et_al_graph_coloring.hpp>
#include <boost/graph/distributed/mpi_process_group.hpp>
#include <boost/lexical_cast.hpp>
#include <boost/graph/parallel/distribution.hpp>
#include <boost/graph/erdos_renyi_generator.hpp>
#include <boost/graph/distributed/adjacency_list.hpp>
#include <boost/graph/graphviz.hpp>
#include <iostream>
#include <boost/random.hpp>
#include <boost/test/minimal.hpp>

#ifdef BOOST_NO_EXCEPTIONS
void
boost::throw_exception(std::exception const& ex)
{
    std::cout << ex.what() << std::endl;
    abort();
}
#endif

using namespace boost;
using boost::graph::distributed::mpi_process_group;

void
test_distributed_graph_coloring(int n, double p, int s,
                                int seed, bool emit_dot_file)
{
  typedef adjacency_list<listS,
                         distributedS<mpi_process_group, vecS>,
                         undirectedS> Graph;

  typedef property_map<Graph, vertex_index_t>::type vertex_index_map;

  // Build a random number generator
  minstd_rand gen;
  gen.seed(seed);

  // Build a random graph
  Graph g(erdos_renyi_iterator<minstd_rand, Graph>(gen, n, p),
          erdos_renyi_iterator<minstd_rand, Graph>(),
          n);

  // Set up color map
  std::vector<int> colors_vec(num_vertices(g));
  iterator_property_map<int*, vertex_index_map> color(&colors_vec[0],
                                                      get(vertex_index, g));

  // Run the graph coloring algorithm
  graph::boman_et_al_graph_coloring(g, color, s);

  if (process_id(g.process_group()) == 0) {
    graph::distributed::boman_et_al_graph_coloring_stats.print(std::cout);
  }

  if ( emit_dot_file ) {
    if ( process_id(g.process_group()) == 0 ) {
      for (int i = 0; i < n; ++i)
        get(color, vertex(i, g));
      synchronize(color);
    } else {
      synchronize(color);
    }

    write_graphviz("coloring.dot", g, paint_by_number(color));
  }
}

int test_main(int argc, char* argv[])
{
  mpi::environment env(argc, argv);

  int n = 1000;
  double p = 0.01;
  int s = 100;
  int seed = 1;
  bool emit_dot_file = false;

  if (argc > 1) n             = lexical_cast<int>(argv[1]);
  if (argc > 2) p             = lexical_cast<double>(argv[2]);
  if (argc > 3) s             = lexical_cast<int>(argv[3]);
  if (argc > 4) seed          = lexical_cast<int>(argv[4]);
  if (argc > 5) emit_dot_file = lexical_cast<bool>(argv[5]);

  test_distributed_graph_coloring(n, p, s, seed, emit_dot_file);

  return 0;
}

max_stars_repo_name: Abce/boost


text:
/*

Copyright (c) 2005-2019, University of Oxford.
All rights reserved.

University of Oxford means the Chancellor, Masters and Scholars of the
University of Oxford, having an administrative office at Wellington
Square, Oxford OX1 2JD, UK.

This file is part of Chaste.

Redistribution and use in source and binary forms, with or without
modification, are permitted provided that the following conditions are met:
 * Redistributions of source code must retain the above copyright notice,
   this list of conditions and the following disclaimer.
 * Redistributions in binary form must reproduce the above copyright notice,
   this list of conditions and the following disclaimer in the documentation
   and/or other materials provided with the distribution.
 * Neither the name of the University of Oxford nor the names of its
   contributors may be used to endorse or promote products derived from this
   software without specific prior written permission.

THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE
GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION)
HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT
OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.

*/

#ifndef CELLIDWRITER_HPP_
#define CELLIDWRITER_HPP_

#include "ChasteSerialization.hpp"
#include <boost/serialization/base_object.hpp>
#include "AbstractCellWriter.hpp"

/**
 * A class written using the visitor pattern for writing cell IDs to file.
 *
 * The output file is called loggedcell.dat by default. If VTK is switched on,
 * then the writer also specifies the VTK output for each cell, which is stored in
 * the VTK cell data "Cell IDs" by default.
 */
template<unsigned ELEMENT_DIM, unsigned SPACE_DIM>
class CellIdWriter : public AbstractCellWriter<ELEMENT_DIM, SPACE_DIM>
{
private:
    /** Needed for serialization. */
    friend class boost::serialization::access;
    /**
     * Serialize the object and its member variables.
     *
     * @param archive the archive
     * @param version the current version of this class
     */
    template<class Archive>
    void serialize(Archive & archive, const unsigned int version)
    {
        archive & boost::serialization::base_object<AbstractCellWriter<ELEMENT_DIM, SPACE_DIM> >(*this);
    }

public:

    /**
     * Default constructor.
     */
    CellIdWriter();

    /**
     * Overridden GetCellDataForVtkOutput() method.
     *
     * Get a double associated with a cell. This method reduces duplication
     * of code between the methods VisitCell() and AddVtkData().
     *
     * @param pCell a cell
     * @param pCellPopulation a pointer to the cell population owning the cell
     *
     * @return data associated with the cell
     */
    double GetCellDataForVtkOutput(CellPtr pCell, AbstractCellPopulation<ELEMENT_DIM, SPACE_DIM>* pCellPopulation);

    /**
     * Overridden VisitCell() method.
     *
     * Visit a cell and write its ID.
     *
     * Outputs a line of space-separated values of the form:
     * ...[cell id] [location index] [x-pos] [y-pos] [z-pos] ...
     * with [y-pos] and [z-pos] included for 2 and 3 dimensional simulations, respectively.
     *
     * This is appended to the output written by AbstractCellBasedWriter, which is a single
     * value [present simulation time], followed by a tab.
     *
     * @param pCell a cell
     * @param pCellPopulation a pointer to the cell population owning the cell
     */
    virtual void VisitCell(CellPtr pCell, AbstractCellPopulation<ELEMENT_DIM, SPACE_DIM>* pCellPopulation);
};

#include "SerializationExportWrapper.hpp"
EXPORT_TEMPLATE_CLASS_ALL_DIMS(CellIdWriter)

#endif /* CELLIDWRITER_HPP_ */

ext: hpp


text: #include <boost/smart_ptr/make_unique.hpp>

max_stars_repo_name: miathedev/BoostForArduino

max_issues_repo_name: miathedev/BoostForArduino

max_forks_repo_name: miathedev/BoostForArduino


text:
#define BOOST_TEST_MODULE SolutionTest

#include "solution.hpp"

//#define BOOST_TEST_DYN_LINK
#include <boost/test/unit_test.hpp>

BOOST_AUTO_TEST_SUITE(SolutionSuite)

BOOST_AUTO_TEST_CASE(PlainTest1)
{
    vector<vector<int>> intervals{{1,3},{2,6},{8,10},{15,18}};
    vector<vector<int>> result = Solution().merge(intervals);

    vector<vector<int>> expected{{1,6},{8,10},{15,18}};

    BOOST_REQUIRE_EQUAL(result.size(), expected.size());
    for (int i = 0; i < result.size(); ++i)
        BOOST_CHECK_EQUAL_COLLECTIONS(result[i].begin(), result[i].end(), expected[i].begin(), expected[i].end());
}

BOOST_AUTO_TEST_CASE(PlainTest2)
{
    vector<vector<int>> intervals{{1,4},{4,5}};
    vector<vector<int>> result = Solution().merge(intervals);

    vector<vector<int>> expected{{1,5}};

    BOOST_REQUIRE_EQUAL(result.size(), expected.size());
    for (int i = 0; i < result.size(); ++i)
        BOOST_CHECK_EQUAL_COLLECTIONS(result[i].begin(), result[i].end(), expected[i].begin(), expected[i].end());
}

BOOST_AUTO_TEST_SUITE_END()

ext: cpp

max_stars_repo_name: johnhany/leetcode

max_issues_repo_name: johnhany/leetcode

max_forks_repo_name: johnhany/leetcode


text:
#ifndef FM_ICLIENT_HPP
#define FM_ICLIENT_HPP

#include "event/input/IInputEvent.hpp"

#include "event/output/FacadeEvent.hpp"

#include "core/CoreClient.hpp"

#include "traffic/socket/ISocket.hpp"

#include "system/IStateMachine.hpp"
#include "system/Timer.hpp"

#include "common/channel_management.pb.h"
#include "common/location.pb.h"

#include <boost/asio.hpp>

#include <memory>
#include <mutex>
#include <functional>

namespace fm
{

namespace state
{
    class IState;
}

namespace backend
{
    class ClientBackend;
}

/**
 * Created by: Bartosz Nawrot
 * Date: 2018-11-25
 * Description:
 */
class IClient : public system::IStateMachine<event::input::IInputEvent>
{
public:
    class Listener
    {
    public:
        virtual ~Listener();

        virtual void onEvent(const std::shared_ptr<const event::output::FacadeEvent>) = 0;

        virtual std::unique_ptr<com::fleetmgr::interfaces::Location> getLocation() = 0;

        virtual void trace(const std::string&) = 0;
    };

    virtual ~IClient();

    void trace(const std::string& message);

    virtual std::string toString() const = 0;

protected:
    std::unique_ptr<backend::ClientBackend> backend;

    IClient(const std::string&, const int, const std::string&, Listener&, boost::asio::io_service&);

private:
    IClient() = delete;
    IClient(IClient&) = delete;

    Listener& listener;
};

} // fm

#endif // FM_IDEVICE_HPP


text:
// Copyright (c) 2019 Duality Blockchain Solutions Developers
// Distributed under the MIT software license, see the accompanying
// file COPYING or http://www.opensource.org/licenses/mit-license.php.

#include "bdap/domainentrydb.h"

#include "amount.h"
#include "base58.h"
#include "bdap/fees.h"
#include "coins.h"
#include "bdap/utils.h"
#include "utilmoneystr.h"
#include "utiltime.h"
#include "validation.h"
#include "validationinterface.h"

#include <univalue.h>

#include <boost/thread.hpp>

CDomainEntryDB *pDomainEntryDB = NULL;

bool GetDomainEntry(const std::vector<unsigned char>& vchObjectPath, CDomainEntry& entry)
{
    if (!pDomainEntryDB || !pDomainEntryDB->ReadDomainEntry(vchObjectPath, entry))
        return false;

    return !entry.IsNull();
}

bool GetDomainEntryPubKey(const std::vector<unsigned char>& vchPubKey, CDomainEntry& entry)
{
    if (!pDomainEntryDB || !pDomainEntryDB->ReadDomainEntryPubKey(vchPubKey, entry))
        return false;

    return !entry.IsNull();
}

bool AccountPubKeyExists(const std::vector<unsigned char>& vchPubKey)
{
    CDomainEntry entry;
    return GetDomainEntryPubKey(vchPubKey, entry);
}

bool DomainEntryExists(const std::vector<unsigned char>& vchObjectPath)
{
    if (!pDomainEntryDB)
        return false;

    return pDomainEntryDB->DomainEntryExists(vchObjectPath);
}

bool DeleteDomainEntry(const CDomainEntry& entry)
{
    if (!pDomainEntryDB)
        return false;

    bool fEraseEntryResult = pDomainEntryDB->EraseDomainEntry(entry.vchFullObjectPath());
    bool fErasePubKeyResult = pDomainEntryDB->EraseDomainEntryPubKey(entry.DHTPublicKey);
    return (fEraseEntryResult && fErasePubKeyResult);
}

bool UndoAddDomainEntry(const CDomainEntry& entry)
{
    if (!pDomainEntryDB)
        return false;

    bool fEraseEntryResult = pDomainEntryDB->EraseDomainEntry(entry.vchFullObjectPath());
    bool fErasePubKeyResult = pDomainEntryDB->EraseDomainEntryPubKey(entry.DHTPublicKey);
    return (fEraseEntryResult && fErasePubKeyResult);
}

bool UndoUpdateDomainEntry(const CDomainEntry& entry)
{
    if (!pDomainEntryDB)
        return false;
    // TODO (BDAP): Implement undo update domain entry used in DisconnectBlock to handle forks
    return false;
}

bool UndoDeleteDomainEntry(const CDomainEntry& entry)
{
    if (!pDomainEntryDB)
        return false;
    // TODO (BDAP): Implement undo delete domain entry used in DisconnectBlock to handle forks
    return false;
}

bool CDomainEntryDB::AddDomainEntry(const CDomainEntry& entry, const int op)
{
    bool writeState = false;
    {
        LOCK(cs_bdap_entry);
        writeState = Write(make_pair(std::string("dc"), entry.vchFullObjectPath()), entry)
                         && Write(make_pair(std::string("pk"), entry.DHTPublicKey), entry);
    }
    if (writeState)
        AddDomainEntryIndex(entry, op);

    return writeState;
}

void CDomainEntryDB::AddDomainEntryIndex(const CDomainEntry& entry, const int op)
{
    UniValue oName(UniValue::VOBJ);
    if (BuildBDAPJson(entry, oName)) {
        CharString vchOperationType = vchFromString(BDAPFromOp(op));
        GetMainSignals().NotifyBDAPUpdate(oName.write().c_str(), reinterpret_cast<char*>(vchOperationType.data()));
        WriteDomainEntryIndexHistory(entry, op);
    }
}

bool CDomainEntryDB::ReadDomainEntry(const std::vector<unsigned char>& vchObjectPath, CDomainEntry& entry)
{
    LOCK(cs_bdap_entry);
    return CDBWrapper::Read(make_pair(std::string("dc"), vchObjectPath), entry);
}

bool CDomainEntryDB::ReadDomainEntryPubKey(const std::vector<unsigned char>& vchPubKey, CDomainEntry& entry)
{
    LOCK(cs_bdap_entry);
    return CDBWrapper::Read(make_pair(std::string("pk"), vchPubKey), entry);
}

bool CDomainEntryDB::EraseDomainEntry(const std::vector<unsigned char>& vchObjectPath)
{
    LOCK(cs_bdap_entry);
    CDomainEntry entry;
    if (!ReadDomainEntry(vchObjectPath, entry)) {
        LogPrintf("CDomainEntryDB::%s -- ReadDomainEntry failed. vchObjectPath = %s\n", __func__, stringFromVch(vchObjectPath));
        return false;
    }

    return CDBWrapper::Erase(make_pair(std::string("dc"), vchObjectPath));
}

bool CDomainEntryDB::EraseDomainEntryPubKey(const std::vector<unsigned char>& vchPubKey)
{
    LOCK(cs_bdap_entry);
    CDomainEntry entry;
    if (!ReadDomainEntryPubKey(vchPubKey, entry))
        return false;

    return CDBWrapper::Erase(make_pair(std::string("pk"), vchPubKey));
}

bool CDomainEntryDB::DomainEntryExists(const std::vector<unsigned char>& vchObjectPath)
{
    LOCK(cs_bdap_entry);
    return CDBWrapper::Exists(make_pair(std::string("dc"), vchObjectPath));
}

bool CDomainEntryDB::DomainEntryExistsPubKey(const std::vector<unsigned char>& vchPubKey)
{
    LOCK(cs_bdap_entry);
    return CDBWrapper::Exists(make_pair(std::string("pk"), vchPubKey));
}

bool CDomainEntryDB::RemoveExpired(int& entriesRemoved)
{
    boost::scoped_ptr<CDBIterator> pcursor(NewIterator());
    pcursor->SeekToFirst();
    CDomainEntry entry;
    std::pair<std::string, std::vector<unsigned char> > key;
    while (pcursor->Valid()) {
        boost::this_thread::interruption_point();
        try {
            if (pcursor->GetKey(key) && key.first == "dc") {
                pcursor->GetValue(entry);
                  if ((unsigned int)chainActive.Tip()->GetMedianTimePast() >= entry.nExpireTime)
                {
                    entriesRemoved++;
                    EraseDomainEntry(key.second);
                }
            }
            else if (pcursor->GetKey(key) && key.first == "txid") {
                std::vector<unsigned char> value;
                CDomainEntry entry;
                pcursor->GetValue(value);
                if (GetDomainEntry(value, entry) && (unsigned int)chainActive.Tip()->GetMedianTimePast() >= entry.nExpireTime)
                {
                    entriesRemoved++;
                    EraseDomainEntryPubKey(entry.DHTPublicKey);
                }
            }
            pcursor->Next();
        } catch (std::exception &e) {
            return error("%s() : deserialize error", __PRETTY_FUNCTION__);
        }
    }
    return true;
}

void CDomainEntryDB::WriteDomainEntryIndexHistory(const CDomainEntry& entry, const int op)
{
    if (IsArgSet("-zmqpubbdaphistory")) {
        UniValue oName(UniValue::VOBJ);
        BuildBDAPJson(entry, oName);
        oName.push_back(Pair("op", BDAPFromOp(op)));
        GetMainSignals().NotifyBDAPUpdate(oName.write().c_str(), "bdap_history");
    }
}

void CDomainEntryDB::WriteDomainEntryIndex(const CDomainEntry& entry, const int op)
{
    if (IsArgSet("-zmqpubbdaprecord")) {
        UniValue oName(UniValue::VOBJ);
        CDynamicAddress address(EncodeBase58(entry.WalletAddress));
        oName.push_back(Pair("address", address.ToString()));
        oName.push_back(Pair("expires_on", entry.nExpireTime));
        oName.push_back(Pair("dht_publickey", HexStr(entry.DHTPublicKey)));
        GetMainSignals().NotifyBDAPUpdate(oName.write().c_str(), "bdap_record");
    }
    WriteDomainEntryIndexHistory(entry, op);
}

bool CDomainEntryDB::UpdateDomainEntry(const std::vector<unsigned char>& vchObjectPath, const CDomainEntry& entry)
{
    LOCK(cs_bdap_entry);

    if (!EraseDomainEntry(vchObjectPath)) {
        LogPrintf("CDomainEntryDB::%s -- EraseDomainEntry failed. vchObjectPath = %s\n", __func__, stringFromVch(vchObjectPath));
        return false;
    }
    if (!EraseDomainEntryPubKey(entry.DHTPublicKey)) {
        LogPrintf("CDomainEntryDB::%s -- EraseDomainEntryPubKey failed. vchObjectPath = %s\n", __func__, stringFromVch(entry.DHTPublicKey));
        return false;
    }

    bool writeState = false;
    writeState = Update(make_pair(std::string("dc"), entry.vchFullObjectPath()), entry)
                    && Update(make_pair(std::string("pk"), entry.DHTPublicKey), entry);
    if (writeState)
        AddDomainEntryIndex(entry, OP_BDAP_MODIFY);

    return writeState;
}

// Removes expired records from databases.
bool CDomainEntryDB::CleanupLevelDB(int& nRemoved)
{
    boost::scoped_ptr<CDBIterator> pcursor(NewIterator());
    pcursor->SeekToFirst();
    CDomainEntry dirEntry;
    std::pair<std::string, std::vector<unsigned char> > key;
    while (pcursor->Valid()) {
        boost::this_thread::interruption_point();
        try {
            if (pcursor->GetKey(key) && key.first == "dc")
            {
                pcursor->GetValue(dirEntry);
                if ((unsigned int)chainActive.Tip()->GetMedianTimePast() >= dirEntry.nExpireTime)
                {
                    nRemoved++;
                    EraseDomainEntry(key.second);
                }
            }
            else if (pcursor->GetKey(key) && key.first == "pk")
            {
                pcursor->GetValue(dirEntry);
                if ((unsigned int)chainActive.Tip()->GetMedianTimePast() >= dirEntry.nExpireTime)
                {
                    nRemoved++;
                    EraseDomainEntryPubKey(dirEntry.DHTPublicKey);
                }
            }
            pcursor->Next();
        } catch (std::exception& e) {
            return error("%s() : deserialize error", __PRETTY_FUNCTION__);
        }
    }
    return true;
}

// Lists active entries by domain name with paging support
bool CDomainEntryDB::ListDirectories(const std::vector<unsigned char>& vchObjectLocation, const unsigned int& nResultsPerPage, const unsigned int& nPage, UniValue& oDomainEntryList, const BDAP::ObjectType& accountType, const std::string searchString)
{
    // TODO: (bdap) implement paging
    // if vchObjectLocation is empty, list entries from all domains
    int index = 0;
    bool addEntry = true;
    std::pair<std::string, CharString> key;
    std::unique_ptr<CDBIterator> pcursor(NewIterator());
    pcursor->SeekToFirst();
    while (pcursor->Valid()) {
        boost::this_thread::interruption_point();
        CDomainEntry entry;
        try {
            if (pcursor->GetKey(key) && key.first == "dc") {
                pcursor->GetValue(entry);
                //filter by accountType, unless DEFAULT
                if ((entry.nObjectType == GetObjectTypeInt(accountType)) ||  (accountType == DEFAULT_ACCOUNT_TYPE)) {
                    if (vchObjectLocation.empty() || entry.vchObjectLocation() == vchObjectLocation)
                    {
                        addEntry = true; //always reset to true

                        if (searchString.size() > 0) {
                            //compare to ObjectID and Common Name
                            std::string compareString(entry.ObjectID.begin(), entry.ObjectID.end());
                            std::string compareCommonString(entry.CommonName.begin(), entry.CommonName.end());
                            std::size_t found = compareString.find(searchString);
                            std::size_t foundCommon = compareCommonString.find(searchString);

                            if ((found==std::string::npos) && (foundCommon==std::string::npos))
                                addEntry = false;
                        }

                        if (addEntry)
                        {
                            UniValue oDomainEntryEntry(UniValue::VOBJ);
                            BuildBDAPJson(entry, oDomainEntryEntry, false);
                            oDomainEntryList.push_back(oDomainEntryEntry);
                            index++;
                        }
                    }
                } //if entry.nObjectType
            }
            pcursor->Next();
        }
        catch (std::exception& e) {
            return error("%s() : deserialize error", __PRETTY_FUNCTION__);
        } //try-catch

    }
    return true;
}

bool CDomainEntryDB::GetDomainEntryInfo(const std::vector<unsigned char>& vchFullObjectPath, UniValue& oDomainEntryInfo)
{
    CDomainEntry entry;
    if (!ReadDomainEntry(vchFullObjectPath, entry)) {
        return false;
    }

    if (!BuildBDAPJson(entry, oDomainEntryInfo, false)) {
        return false;
    }

    return true;
}

bool CDomainEntryDB::GetDomainEntryInfo(const std::vector<unsigned char>& vchFullObjectPath, CDomainEntry& entry)
{
    if (!ReadDomainEntry(vchFullObjectPath, entry)) {
        return false;
    }

    return true;
}

bool CheckDomainEntryDB()
{
    if (!pDomainEntryDB)
        return false;

    return true;
}

bool FlushLevelDB()
{
    {
        LOCK(cs_bdap_entry);
        if (pDomainEntryDB != NULL)
        {
            if (!pDomainEntryDB->Flush()) {
                LogPrintf("Failed to write to BDAP database!");
                return false;
            }
        }
    }
    return true;
}

void CleanupLevelDB(int& nRemoved)
{
    if(pDomainEntryDB != NULL)
        pDomainEntryDB->CleanupLevelDB(nRemoved);
    FlushLevelDB();
}

static bool CommonDataCheck(const CDomainEntry& entry, const vchCharString& vvchOpParameters, std::string& errorMessage)
{
    if (entry.IsNull() == true)
    {
        errorMessage = "CommonDataCheck failed! DomainEntry is null.";
        return false;
    }

    if (vvchOpParameters.size() == 0)
    {
        errorMessage = "CommonDataCheck failed! Invalid parameters.";
        return false;
    }

    if (vvchOpParameters.size() != 3)
    {
        errorMessage = "CommonDataCheck failed! Not enough parameters.";
        return false;
    }

    if (entry.GetFullObjectPath() != stringFromVch(vvchOpParameters[0]))
    {
        errorMessage = "CommonDataCheck failed! Script operation parameter does not match entry entry object.";
        return false;
    }

    if (entry.DHTPublicKey != vvchOpParameters[1])
    {
        errorMessage = "CommonDataCheck failed! DHT public key mismatch.";
        return false;
    }

    if (vvchOpParameters[2].size() > 10)
    {
        errorMessage = "CommonDataCheck failed! Expire date or expire months invalid.";
        return false;
    }

    if (entry.DomainComponent != vchDefaultDomainName)
    {
        errorMessage = "CommonDataCheck failed! Must use default domain.";
        return false;
    }

    if (entry.OrganizationalUnit == vchDefaultAdminOU)
    {
        errorMessage = "CommonDataCheck failed! Can not use default admin domain.";
        return false;
    }

    if (entry.OrganizationalUnit != vchDefaultPublicOU)
    {
        errorMessage = "CommonDataCheck failed! Must use default public organizational unit.";
        return false;
    }

    return true;
}

static bool CheckNewDomainEntryTxInputs(const CDomainEntry& entry, const CScript& scriptOp, const vchCharString& vvchOpParameters, const uint256& txHash,
                               std::string& errorMessage, bool fJustCheck)
{
    if (!CommonDataCheck(entry, vvchOpParameters, errorMessage))
        return error(errorMessage.c_str());

    if (fJustCheck)
        return true;

    CDomainEntry getDomainEntry;
    if (GetDomainEntry(entry.vchFullObjectPath(), getDomainEntry))
    {
        if (entry.txHash != txHash) {
            errorMessage = "CheckNewDomainEntryTxInputs: - The entry " + getDomainEntry.GetFullObjectPath() + " already exists.  Add new entry failed!";
            return error(errorMessage.c_str());
        }
        else {
            LogPrintf("%s -- Already have entry %s in local database. Skipping add entry step.\n", __func__, entry.GetFullObjectPath());
            return true;
        }
    }

    if (!pDomainEntryDB)
    {
        errorMessage = "CheckNewDomainEntryTxInputs failed! Can not open LevelDB BDAP entry database.";
        return error(errorMessage.c_str());
    }
    int op = OP_BDAP_NEW;
    if (!pDomainEntryDB->AddDomainEntry(entry, op))
    {
        errorMessage = "CheckNewDomainEntryTxInputs failed! Error adding new entry entry request to LevelDB.";
        return error(errorMessage.c_str());
    }

    return FlushLevelDB();
}

static bool CheckDeleteDomainEntryTxInputs(const CDomainEntry& entry, const CScript& scriptOp, const vchCharString& vvchOpParameters,
                                  std::string& errorMessage, bool fJustCheck)
{
    if (vvchOpParameters.size() == 0) {
        errorMessage = "CheckDeleteDomainEntryTxInputs: - Invalid delete operation parameters. This delete operation failed!";
        return error(errorMessage.c_str());
    }
    if (fJustCheck)
        return true;

    std::vector<unsigned char> vchFullObjectPath = vvchOpParameters[0];
    CDomainEntry prevDomainEntry;
    if (!GetDomainEntry(vchFullObjectPath, prevDomainEntry))
    {
        errorMessage = "CheckDeleteDomainEntryTxInputs: - Can not find " + prevDomainEntry.GetFullObjectPath() + " entry; this delete operation failed!";
        return error(errorMessage.c_str());
    }

    CTxDestination bdapDest;
    if (!ExtractDestination(scriptOp, bdapDest))
    {
        errorMessage = "CheckDeleteDomainEntryTxInputs: - " + _("Cannot extract destination of BDAP input; this delete operation failed!");
        return error(errorMessage.c_str());
    }
    else
    {
        CTransactionRef prevTx = MakeTransactionRef();
        uint256 hashBlock;
        if (!GetTransaction(prevDomainEntry.txHash, prevTx, Params().GetConsensus(), hashBlock, true)) {
            errorMessage = "CheckDeleteDomainEntryTxInputs: - " + _("Cannot extract previous transaction from BDAP output; this delete operation failed!");
            return error(errorMessage.c_str());
        }
        // Get current wallet address used for BDAP tx
        CDynamicAddress txAddress = GetScriptAddress(scriptOp);
        // Get previous wallet address used for BDAP tx
        CScript prevScriptPubKey;
        GetBDAPOpScript(prevTx, prevScriptPubKey);
        CDynamicAddress prevAddress = GetScriptAddress(prevScriptPubKey);
        if (txAddress.ToString() != prevAddress.ToString())
        {
            //check if previous wallet address is used for update and delete txs
            errorMessage = "CheckDeleteDomainEntryTxInputs: - " + _("Delete must use the previous wallet address; this delete operation failed!");
            return error(errorMessage.c_str());
        }
    }

    //Remove PubKey entry also (2 records in LevelDB for each domainentry)
    pDomainEntryDB->EraseDomainEntryPubKey(prevDomainEntry.DHTPublicKey);

    if (!pDomainEntryDB->EraseDomainEntry(vchFullObjectPath))
    {
        errorMessage = "CheckDeleteDomainEntryTxInputs: - Error deleting entry entry in LevelDB; this delete operation failed!";
        return error(errorMessage.c_str());
    }

    return FlushLevelDB();
}

static bool CheckUpdateDomainEntryTxInputs(CDomainEntry& entry, const CScript& scriptOp, const vchCharString& vvchOpParameters, const uint256& txHash, const int& nMonths, const uint32_t& nBlockTime,
                                  std::string& errorMessage, bool fJustCheck)
{
    //if exists, check for owner's signature
    if (!CommonDataCheck(entry, vvchOpParameters, errorMessage))
        return error(errorMessage.c_str());

    if (fJustCheck)
        return true;

    CDomainEntry prevDomainEntry;
    if (!GetDomainEntry(entry.vchFullObjectPath(), prevDomainEntry))
    {
        if (entry.txHash != txHash) {
            errorMessage = "CheckUpdateDomainEntryTxInputs: - Can not find " + prevDomainEntry.GetFullObjectPath() + " entry; this update operation failed!";
            return error(errorMessage.c_str());
        }
        else {
            LogPrintf("%s -- Already have entry %s in local database. Skipping update entry step.\n", __func__, entry.GetFullObjectPath());
            return true;
        }
    }

    CTxDestination bdapDest;
    if (!ExtractDestination(scriptOp, bdapDest))
    {
        errorMessage = "CheckUpdateDomainEntryTxInputs: - " + _("Cannot extract destination of BDAP input; this update operation failed!");
        return error(errorMessage.c_str());
    }
    else
    {
        CTransactionRef prevTx = MakeTransactionRef();
        uint256 hashBlock;
        if (!GetTransaction(prevDomainEntry.txHash, prevTx, Params().GetConsensus(), hashBlock, true)) {
            errorMessage = "CheckUpdateDomainEntryTxInputs: - " + _("Cannot extract previous transaction from BDAP output; this update operation failed!");
            return error(errorMessage.c_str());
        }
        // Get current wallet address used for BDAP tx
        CDynamicAddress txAddress = GetScriptAddress(scriptOp);
        // Get previous wallet address used for BDAP tx
        CScript prevScriptPubKey;
        GetBDAPOpScript(prevTx, prevScriptPubKey);
        CDynamicAddress prevAddress = GetScriptAddress(prevScriptPubKey);
        if (txAddress.ToString() != prevAddress.ToString())
        {
            //check if previous wallet address is used for update and delete txs
            errorMessage = "CheckUpdateDomainEntryTxInputs: - " + _("Update must use the previous wallet address; this update operation failed!");
            return error(errorMessage.c_str());
        }
    }
    entry.nExpireTime = AddMonthsToBlockTime(prevDomainEntry.nExpireTime, nMonths);
    LogPrint("bdap", "%s -- prevDomainEntry.nExpireTime = %d, AddMonthsToBlockTime() = %d, nMonths = %d\n", __func__,
                    prevDomainEntry.nExpireTime, AddMonthsToBlockTime(prevDomainEntry.nExpireTime, nMonths), nMonths);
    if (!pDomainEntryDB->UpdateDomainEntry(entry.vchFullObjectPath(), entry))
    {
        errorMessage = "CheckUpdateDomainEntryTxInputs: - Error updating entry in LevelDB; this update operation failed!";
        return error(errorMessage.c_str());
    }

    return FlushLevelDB();
}

static bool CheckMoveDomainEntryTxInputs(const CDomainEntry& entry, const CScript& scriptOp, const vchCharString& vvchOpParameters,
                                std::string& errorMessage, bool fJustCheck)
{
    //check name in operation matches entry data in leveldb
    //check if exists already
    //if exists, check for owner's signature
    return false;
}

bool CheckDomainEntryTx(const CTransactionRef& tx, const CScript& scriptOp, const int& op1, const int& op2, const std::vector<std::vector<unsigned char> >& vvchArgs,
                                const bool fJustCheck, const int& nHeight, const uint32_t& nBlockTime, const bool bSanityCheck, std::string& errorMessage)
{
    if (tx->IsCoinBase() && !fJustCheck && !bSanityCheck)
    {
        LogPrintf("*Trying to add BDAP entry in coinbase transaction, skipping...");
        return true;
    }

    LogPrint("bdap", "%s -- BDAP nHeight=%d, chainActive.Tip()=%d, op1=%s, op2=%s, hash=%s justcheck=%s\n", __func__, nHeight, chainActive.Tip()->nHeight, BDAPFromOp(op1).c_str(), BDAPFromOp(op2).c_str(), tx->GetHash().ToString().c_str(), fJustCheck ? "JUSTCHECK" : "BLOCK");

    // unserialize BDAP from txn, check if the entry is valid and does not conflict with a previous entry
    CDomainEntry entry;
    std::vector<unsigned char> vchData;
    std::vector<unsigned char> vchHash;
    int nDataOut;
    const std::string strOperationType = GetBDAPOpTypeString(op1, op2);
    if (strOperationType != "bdap_delete_account") {
        bool bData = GetBDAPData(tx, vchData, vchHash, nDataOut);
        if(bData && !entry.UnserializeFromData(vchData, vchHash))
        {
            errorMessage = "BDAP_CONSENSUS_ERROR: ERRCODE: 3601 - " + _("UnserializeFromData data in tx failed!");
            LogPrintf("%s -- %s \n", __func__, errorMessage);
            return error(errorMessage.c_str());
        }

        if(!entry.ValidateValues(errorMessage))
        {
            errorMessage = "BDAP_CONSENSUS_ERROR: ERRCODE: 3602 - " + errorMessage;
            LogPrintf("%s -- %s \n", __func__, errorMessage);
            return error(errorMessage.c_str());
        }

        entry.txHash = tx->GetHash();
        entry.nHeight = nHeight;
    }

    CAmount monthlyFee, oneTimeFee, depositFee;
    if (strOperationType == "bdap_new_account") {
        if (vvchArgs.size() != 3) {
            errorMessage = "Failed to get fees to add a new BDAP account";
            return false;
        }
        std::string strMonths = stringFromVch(vvchArgs[2]);
        std::size_t foundMonth = strMonths.find("Month");
        if (foundMonth != std::string::npos)
            strMonths.replace(foundMonth, 5, "");

        uint32_t nMonths;
        ParseUInt32(strMonths, &nMonths);
        if (nMonths > MAX_REGISTRATION_MONTHS)
            nMonths = MAX_REGISTRATION_MONTHS;

        if (!GetBDAPFees(OP_BDAP_NEW, OP_BDAP_ACCOUNT_ENTRY, entry.ObjectType(), (uint16_t)nMonths, monthlyFee, oneTimeFee, depositFee)) {
            errorMessage = "Failed to get fees to add a new BDAP account";
            return false;
        }
        LogPrint("bdap", "%s -- nMonths %d, monthlyFee %d, oneTimeFee %d, depositFee %d\n", __func__,
                                nMonths, FormatMoney(monthlyFee), FormatMoney(oneTimeFee), FormatMoney(depositFee));
        // extract amounts from tx.
        CAmount dataAmount, opAmount;
        if (!ExtractAmountsFromTx(tx, dataAmount, opAmount)) {
            errorMessage = "Unable to extract BDAP amounts from transaction";
            return false;
        }
        LogPrint("bdap", "%s -- dataAmount %d, opAmount %d\n", __func__, FormatMoney(dataAmount), FormatMoney(opAmount));
        if (monthlyFee > dataAmount) {
            LogPrintf("%s -- Invalid BDAP monthly registration fee amount for new BDAP account. Monthly paid %d but should be %d\n", __func__,
                                    FormatMoney(dataAmount), FormatMoney(monthlyFee));
            errorMessage = "Invalid BDAP monthly registration fee amount for new BDAP account";
            return false;
        }
        else {
            LogPrint("bdap", "%s -- Valid BDAP monthly registration fee amount for new BDAP account. Monthly paid %d, should be %d.\n", __func__,
                                    FormatMoney(dataAmount), FormatMoney(monthlyFee));
        }
        if (depositFee > opAmount) {
            LogPrintf("%s -- Invalid BDAP deposit fee amount for new BDAP account. Deposit paid %d but should be %d\n", __func__,
                                    FormatMoney(opAmount), FormatMoney(depositFee));
            errorMessage = "Invalid BDAP deposit fee amount for new BDAP account";
            return false;
        }
        else {
            LogPrint("bdap", "%s -- Valid BDAP deposit fee amount for new BDAP account. Deposit paid %d, should be %d\n", __func__,
                                    FormatMoney(opAmount), FormatMoney(depositFee));
        }
        entry.nExpireTime = AddMonthsToBlockTime(nBlockTime, nMonths);

        return CheckNewDomainEntryTxInputs(entry, scriptOp, vvchArgs, tx->GetHash(), errorMessage, fJustCheck);
    }
    else if (strOperationType == "bdap_delete_account") {
        uint16_t nMonths = 0;
        if (!GetBDAPFees(OP_BDAP_DELETE, OP_BDAP_ACCOUNT_ENTRY, entry.ObjectType(), nMonths, monthlyFee, oneTimeFee, depositFee)) {
            errorMessage = "Failed to get fees to delete a BDAP account";
            return false;
        }

        return CheckDeleteDomainEntryTxInputs(entry, scriptOp, vvchArgs, errorMessage, fJustCheck);
    }
    else if (strOperationType == "bdap_update_account") {
        if (vvchArgs.size() != 3) {
            errorMessage = "Failed to get fees to add a new BDAP account";
            return false;
        }
        std::string strMonths = stringFromVch(vvchArgs[2]);
        std::size_t foundMonth = strMonths.find("Month");
        if (foundMonth != std::string::npos)
            strMonths.replace(foundMonth, 5, "");

        uint32_t nMonths;
        ParseUInt32(strMonths, &nMonths);
        if (nMonths > MAX_REGISTRATION_MONTHS)
            nMonths = MAX_REGISTRATION_MONTHS;

        if (!GetBDAPFees(OP_BDAP_MODIFY, OP_BDAP_ACCOUNT_ENTRY, entry.ObjectType(), (uint16_t)nMonths, monthlyFee, oneTimeFee, depositFee)) {
            errorMessage = "Failed to get fees to add a new BDAP account";
            return false;
        }

        LogPrint("bdap", "%s -- nMonths %d, monthlyFee %d, oneTimeFee %d, depositFee %d\n", __func__,
                                nMonths, FormatMoney(monthlyFee), FormatMoney(oneTimeFee), FormatMoney(depositFee));
        // extract amounts from tx.
        CAmount dataAmount, opAmount;
        if (!ExtractAmountsFromTx(tx, dataAmount, opAmount)) {
            errorMessage = "Unable to extract BDAP amounts from transaction";
            return false;
        }
        LogPrint("bdap", "%s -- dataAmount %d, opAmount %d\n", __func__, FormatMoney(dataAmount), FormatMoney(opAmount));
        if (monthlyFee > dataAmount) {
            LogPrintf("%s -- Invalid BDAP data fee amount for updated BDAP account. Total paid %d but should be %d\n", __func__,
                                dataAmount, monthlyFee);
            errorMessage = "Invalid BDAP deposit fee amount for updated BDAP account";
            return false;
        } else {
            LogPrint("bdap", "%s -- Valid BDAP data fee amount for updated BDAP account. Total paid %d, should be %d\n", __func__,
                                FormatMoney(dataAmount), FormatMoney(monthlyFee));
        }
        if (oneTimeFee > opAmount) {
            LogPrintf("%s -- Invalid BDAP one-time fee amount for updated BDAP account. Total paid %d but should be %d\n", __func__,
                                FormatMoney(dataAmount), FormatMoney(monthlyFee));
            errorMessage = "Invalid BDAP one-time fee amount for updated BDAP account";
            return false;
        } else {
            LogPrint("bdap", "%s -- Valid BDAP one-time fee amount for updated BDAP account. Total paid %d, should be %d\n", __func__,
                                FormatMoney(dataAmount), FormatMoney(monthlyFee));
        }
        // Add previous expire date plus additional months
        return CheckUpdateDomainEntryTxInputs(entry, scriptOp, vvchArgs, tx->GetHash(), nMonths, nBlockTime, errorMessage, fJustCheck);
    }
    else if (strOperationType == "bdap_move_account") {
        uint16_t nMonths = 0;
        if (!GetBDAPFees(OP_BDAP_MOVE, OP_BDAP_ACCOUNT_ENTRY, entry.ObjectType(), nMonths, monthlyFee, oneTimeFee, depositFee)) {
            errorMessage = "Failed to get fees to move a BDAP account to another domain";
            return false;
        }

        return CheckMoveDomainEntryTxInputs(entry, scriptOp, vvchArgs, errorMessage, fJustCheck);
    }

    return false;
}

max_stars_repo_name: blackrangersoftware/brs-dynamic-core

max_issues_repo_name: yangbaitu/Dynamic

max_forks_repo_name: yangbaitu/Dynamic


text:
////////////////////////////////////////////////////////////////////////////////
/// DISCLAIMER
///
/// Copyright 2014-2020 ArangoDB GmbH, Cologne, Germany
/// Copyright 2004-2014 triAGENS GmbH, Cologne, Germany
///
/// Licensed under the Apache License, Version 2.0 (the "License");
/// you may not use this file except in compliance with the License.
/// You may obtain a copy of the License at
///
///     http://www.apache.org/licenses/LICENSE-2.0
///
/// Unless required by applicable law or agreed to in writing, software
/// distributed under the License is distributed on an "AS IS" BASIS,
/// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
/// See the License for the specific language governing permissions and
/// limitations under the License.
///
/// Copyright holder is ArangoDB GmbH, Cologne, Germany
///
/// @author Dan Larkin-York
/// @author Copyright 2017, ArangoDB GmbH, Cologne, Germany
////////////////////////////////////////////////////////////////////////////////

#include <memory>

#include <boost/asio/io_service.hpp>

#undef sleep
#include <boost/thread/thread.hpp>

#include "MockScheduler.h"

using namespace arangodb::cache;

MockScheduler::MockScheduler(std::size_t threads)
    : _ioService(new asio_ns::io_context()),
      _serviceGuard(new asio_ns::io_context::work(*_ioService)) {
  for (std::size_t i = 0; i < threads; i++) {
    auto worker = std::bind(static_cast<size_t (asio_ns::io_context::*)()>(
                                &asio_ns::io_context::run),
                            _ioService.get());
    _group.emplace_back(new std::thread(worker));
  }
}

MockScheduler::~MockScheduler() {
  _serviceGuard.reset();
  for (auto g : _group) {
    g->join();
    delete g;
  }
  _ioService->stop();
}

void MockScheduler::post(std::function<void()> fn) { _ioService->post(fn); }

max_stars_repo_name: Mu-L/arangodb

max_issues_repo_name: Mu-L/arangodb

max_forks_repo_name: Mu-L/arangodb


text:
/*******************************************************************************
*
* MIT License
*
* Copyright (c) 2020 Advanced Micro Devices, Inc.
*
* Permission is hereby granted, free of charge, to any person obtaining a copy
* of this software and associated documentation files (the "Software"), to deal
* in the Software without restriction, including without limitation the rights
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
* copies of the Software, and to permit persons to whom the Software is
* furnished to do so, subject to the following conditions:
*
* The above copyright notice and this permission notice shall be included in all
* copies or substantial portions of the Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
* SOFTWARE.
*
*******************************************************************************/

#include <miopen/solver.hpp>
#include <miopen/env.hpp>
#include <miopen/stringutils.hpp>
#include <miopen/sequences.hpp>
#include <miopen/kernel_build_params.hpp>
#include <miopen/generic_search.hpp>
#include <miopen/conv/data_invoke_params.hpp>
#include <miopen/conv/compiled_in_parameters.hpp>
#include <miopen/conv/wrw_invoke_params.hpp>

#include <boost/any.hpp>

#include <tuple>

MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3)
MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3_PERF_VALS)

#define WINODATA 2
#define WINOFILTER 3
#define MAX_CU_LIMIT 512

static inline size_t Ceil(const size_t v, const size_t m)
{
    assert(m > 0);
    return (v + m - 1) / m;
}

static inline size_t quantize_up(size_t val, size_t factor) { return Ceil(val, factor) * factor; }

static inline int GetBestNGroupParam(const int R,
                                     const int S,
                                     const int R_stride,
                                     const int S_stride,
                                     const int C,
                                     const int K,
                                     const int OH,
                                     const int OW,
                                     const int pad_H,
                                     const int pad_W,
                                     const int N,
                                     const int idilation_w,
                                     const int idilation_h,
                                     const int n_groups,
                                     const int G)
{
    int o_tile     = WINODATA;
    int f_tile     = WINOFILTER;
    int r_factor   = f_tile * 2;
    int s_factor   = r_factor;
    int c_factor   = 2;
    int k_factor   = 32;
    int nwh_factor = 32;
    int w_factor   = o_tile * idilation_w * S_stride;
    int h_factor   = o_tile * idilation_h * R_stride;

    if(S_stride == 1 && idilation_w == 1 && S <= f_tile)
        s_factor = f_tile;
    if((R_stride == 1 && idilation_h == 1) || (R % (f_tile * 2)) == 1)
        r_factor = f_tile;
    if(S_stride == 2 || R_stride == 2 || idilation_w == 2 || idilation_h == 2)
        c_factor = 1;

    size_t g_s = quantize_up(S, s_factor);
    size_t g_r = quantize_up(R, r_factor);
    size_t g_c = quantize_up(C, c_factor);
    size_t g_k = quantize_up(K, k_factor);
    size_t g_w = OW;
    size_t g_h = OH;

    if((pad_W % 2 == 0) && (idilation_w > 1 || S_stride > 1))
        g_w += 1;
    if((pad_H % 2 == 1) && (idilation_h > 1 || R_stride > 1))
        g_h += 1;

    g_w            = quantize_up(g_w, w_factor);
    g_h            = quantize_up(g_h, h_factor);
    size_t g_n_w_h = quantize_up(g_w * g_h * N, nwh_factor * w_factor * h_factor);

    int best_n_groups_cnt = 1;
    double min_param      = 0;
    for(auto i = 1; i < n_groups; ++i)
    {
        size_t g_n_w_h_k =
            quantize_up(g_n_w_h * g_k, nwh_factor * w_factor * h_factor * k_factor * i);
        size_t granulated_mac_count = g_n_w_h_k * g_c * g_s * g_r;
        size_t n_groups_per_cu      = Ceil(i * G, n_groups);
        double perf_metric = static_cast<double>(n_groups_per_cu) * granulated_mac_count / i;
        if(static_cast<double>(granulated_mac_count) / i > 1.0e+7)
            perf_metric *= (1 + i * 0.003);
        else
            perf_metric *= (1 + i * 0.04);
        if(i == 1)
            min_param = perf_metric;
        if(min_param > perf_metric)
        {
            best_n_groups_cnt = i;
            min_param         = perf_metric;
        }
    }
    return best_n_groups_cnt;
}

namespace miopen {
namespace solver {

namespace {
// clang-format off
    auto PerfFieldRules()
    {
        return seq::MakeRuleSet(
            std::make_tuple(seq::Span<int, 1, MAX_CU_LIMIT>{}, &PerformanceConfigConvBinWinogradRxSf2x3::n_groups)
        );
    }
// clang-format on

/// \todo Consider re-using code from RxS.
inline bool IsShaderContraintsMet(const int R,
                                  const int S,
                                  const int,
                                  const int,
                                  const int C,
                                  const int K,
                                  const int H,
                                  const int W,
                                  const int OH,
                                  const int OW,
                                  const int N,
                                  const ConvolutionContext& params)
{
    // Padding for bwd data shall not be negative.
    /// \todo Either remove WrW related code or re-use function from RxS
    if(params.direction.IsBackwardData() || params.direction.IsBackwardWrW())
    {
        if(!(0 <= params.GetBackwardPadW() && params.GetBackwardPadW() < std::pow(2, 16)))
            return false;
        if(!(0 <= params.GetBackwardPadH() && params.GetBackwardPadH() < std::pow(2, 16)))
            return false;
    }
    const auto grid_workgroup_count_x = params.GetStream().GetMaxComputeUnits();
    assert(params.weights_layout.length() == 0);
    // clang-format off
        // Check implementation limits.
        return N < std::pow(2, 16)
            && C < std::pow(2, 16)
            && K < std::pow(2, 16)
            && H < std::pow(2, 16)
            && W < std::pow(2, 16)
            && OH < std::pow(2, 16)
            && OW < std::pow(2, 16)
            && params.pad_w < std::pow(2, 16)
            && params.pad_h < std::pow(2, 16)
            && S < std::pow(2, 16)
            && R < std::pow(2, 16)
            && grid_workgroup_count_x < std::pow(2, 16)
            && (C * H * W) <= std::pow(2, 28)
            && (OH * OW) <= std::pow(2, 23)
            && (K * OH * OW) <= std::pow(2, 28)
            && (K * R * S) <= std::pow(2, 28)
            && (C * R * S) <= std::pow(2, 28);
    // clang-format on
}

} // namespace

PerformanceConfigConvBinWinogradRxSf2x3::PerformanceConfigConvBinWinogradRxSf2x3(int n_groups_)
    : n_groups(n_groups_)
{
}

void PerformanceConfigConvBinWinogradRxSf2x3::EuristicInit(const ConvolutionContext& config)
{
    const auto n_inputs_per_group  = config.n_inputs / config.group_counts,
               n_outputs_per_group = config.n_outputs / config.group_counts;
    if(config.group_counts == 1)
    {
        n_groups = config.GetStream().GetMaxComputeUnits();
        return;
    }

    if(config.direction.IsBackwardWrW())
    {
        n_groups = GetBestNGroupParam(config.in_height,
                                      config.in_width,
                                      config.kernel_dilation_h,
                                      config.kernel_dilation_w,
                                      config.batch_sz,    // N
                                      n_inputs_per_group, // K
                                      config.kernel_size_h,
                                      config.kernel_size_w,
                                      config.pad_w,
                                      config.pad_h,
                                      n_outputs_per_group, // C
                                      config.kernel_stride_h,
                                      config.kernel_stride_w,
                                      config.GetStream().GetMaxComputeUnits(),
                                      config.group_counts);
    }
    else
    {
        n_groups = GetBestNGroupParam(config.kernel_size_h, // RxS
                                      config.kernel_size_w,
                                      config.kernel_stride_h,
                                      config.kernel_stride_w,
                                      n_inputs_per_group,  // C
                                      n_outputs_per_group, // K
                                      config.out_height,   // OHxOW
                                      config.out_width,
                                      config.pad_w,
                                      config.pad_h,
                                      config.batch_sz, // N
                                      config.kernel_dilation_h,
                                      config.kernel_dilation_w,
                                      config.GetStream().GetMaxComputeUnits(),
                                      config.group_counts);
    }
}

bool PerformanceConfigConvBinWinogradRxSf2x3::SetNextValue()
{
    return !PerfFieldRules().Next(*this);
}

bool PerformanceConfigConvBinWinogradRxSf2x3::IsValidValue() const
{
    return PerfFieldRules().IsIn(*this);
}

bool PerformanceConfigConvBinWinogradRxSf2x3::IsValid(const ConvolutionContext& config) const
{
    if(config.GetStream().GetMaxComputeUnits() < n_groups)
        return false;

    if(!IsValidValue())
        return false;
    return true;
}

inline bool PerformanceConfigConvBinWinogradRxSf2x3::
operator==(const PerformanceConfigConvBinWinogradRxSf2x3& other) const
{
    return n_groups == other.n_groups;
}

std::string PerformanceConfigConvBinWinogradRxSf2x3::ToString() const
{
    std::ostringstream ss;
    Serialize(ss);
    return ss.str();
}

PerformanceConfigConvBinWinogradRxSf2x3
ConvBinWinogradRxSf2x3::GetPerformanceConfig(const ConvolutionContext& params) const
{
    PerformanceConfigConvBinWinogradRxSf2x3 pp;
    pp.EuristicInit(params);
    MIOPEN_LOG_I(pp.ToString());
    return pp;
}

bool ConvBinWinogradRxSf2x3::IsValidPerformanceConfig(
    const ConvolutionContext& problem, const PerformanceConfigConvBinWinogradRxSf2x3& c) const
{
    return c.IsValidValue() && c.IsValid(problem);
}

PerformanceConfigConvBinWinogradRxSf2x3
ConvBinWinogradRxSf2x3::Search(const ConvolutionContext& context) const
{
    if(context.direction.IsForward())
        return GenericSearchFwd(*this, context);
    else
        return GenericSearchBwd(*this, context);
}

inline void FillVarsFromConfig(int& H,
                               int& W,
                               int& R,
                               int& S,
                               int& R_stride,
                               int& S_stride,
                               int& C,
                               int& K,
                               int& out_H,
                               int& out_W,
                               int& pad_H,
                               int& pad_W,
                               int& N,
                               int& idilation_w,
                               int& idilation_h,
                               int& n_groups,
                               int& group_cnt,
                               const ConvolutionContext& config)
{
    group_cnt   = config.group_counts;
    n_groups    = config.GetStream().GetMaxComputeUnits();
    pad_H       = config.direction.IsForward() ? config.pad_h : config.GetBackwardPadH();
    pad_W       = config.direction.IsForward() ? config.pad_w : config.GetBackwardPadW();
    H           = config.in_height;
    W           = config.in_width;
    R           = config.kernel_size_h;
    S           = config.kernel_size_w;
    R_stride    = config.kernel_stride_h;
    S_stride    = config.kernel_stride_w;
    C           = config.n_inputs;
    K           = config.n_outputs;
    out_H       = config.out_height;
    out_W       = config.out_width;
    N           = config.batch_sz;
    idilation_w = config.kernel_dilation_h;
    idilation_h = config.kernel_dilation_w;
}

template <typename B, typename T, typename TW>
int ConvBinWinogradRxSf2x3::RunAndMeasureSolution(const miopen::Handle& profile_h,
                                                  B bot_ocl_buf,
                                                  T top_ocl_buf,
                                                  TW wei_ocl_buf,
                                                  ConstData_t bias_ocl_buf,
                                                  const ConvolutionContext& config,
                                                  const ConvSolution& solution,
                                                  float& elapsed_time) const
{
    assert(bias_ocl_buf == nullptr);
    (void)bias_ocl_buf;
    const KernelInfo k_info = solution.construction_params.back();
#ifdef NDEBUG
    try
#endif
    {
        elapsed_time = std::numeric_limits<float>::max();
        // ConvolutionContext::general_compile_options is for OpenCL kernels
        // and thus not applicable for assembly.
        auto kernel = profile_h.AddKernel("",
                                          "",
                                          k_info.kernel_file,
                                          k_info.kernel_name,
                                          k_info.l_wk,
                                          k_info.g_wk,
                                          k_info.comp_options);

        int reserved      = 0;
        int* reserved_ptr = nullptr;

        static const int F_REVERSE_R     = 1 << 0;
        static const int F_REVERSE_S     = 1 << 1;
        static const int F_FLIP_K_C      = 1 << 2;
        static const int F_NKC_STRIDES   = 1 << 9;
        static const int F_GROUP_STRIDES = 1 << 10;

        const auto is_forward = config.direction.IsForward();
        const auto is_wrw     = config.direction.IsBackwardWrW();

        int H, W, R, S, R_stride, S_stride, C, K, out_H, out_W, pad_H, pad_W, N, idilation_w,
            idilation_h, n_groups, group_cnt;
        // clang-format off
        FillVarsFromConfig(H, W, R, S, R_stride, S_stride, C, K, out_H, out_W, pad_H,
                           pad_W, N, idilation_w, idilation_h, n_groups, group_cnt, config);
        // clang-format on
        int flags = ((is_forward || is_wrw) ? 0 : F_REVERSE_R + F_REVERSE_S + F_FLIP_K_C) +
                    F_NKC_STRIDES + F_GROUP_STRIDES;

        C = C / group_cnt;
        K = K / group_cnt;
        // cppcheck-suppress unreadVariable

        const auto in_layout = (is_wrw) ? GetSwappedNCLayout(GetMemLayout_t(config.in_layout))
                                        : GetMemLayout_t(config.in_layout);
        const auto out_layout = (is_wrw) ? GetSwappedNCLayout(GetMemLayout_t(config.out_layout))
                                         : GetMemLayout_t(config.out_layout);
        const auto wei_layout =
            (is_wrw) ? GetSwappedNCLayout(MemLayout_t::NCHW)
                     : (is_forward ? (MemLayout_t::NCHW) : GetSwappedNCLayout(MemLayout_t::NCHW));

        // cppcheck-suppress unreadVariable
        BuffInfo d_buf(GetGroupConvLayout(in_layout, true),
                       N,
                       C,
                       H,
                       W,
                       group_cnt,
                       GetTypeSize(config.in_data_type)),
            // cppcheck-suppress unreadVariable
            o_buf(GetGroupConvLayout(out_layout, true),
                  N,
                  K,
                  out_H,
                  out_W,
                  group_cnt,
                  GetTypeSize(config.out_data_type)),
            // cppcheck-suppress unreadVariable
            f_buf(GetGroupConvLayout(wei_layout, false),
                  K,
                  C,
                  R,
                  S,
                  group_cnt,
                  GetTypeSize(config.weights_data_type));

        if(k_info.l_wk[0] != 0)
            n_groups = solver::ConvBinWinogradRxSf2x3::GetNGroups(config.group_counts,
                                                                  k_info.g_wk[0] / k_info.l_wk[0]);
        else
            n_groups = solver::ConvBinWinogradRxSf2x3::GetNGroups(
                config.group_counts,
                k_info.g_wk[0] / 512); // For OCL runtime. Issue #1724

        kernel(N,
               C,
               H,
               W,
               K,
               n_groups,
               flags,
               reserved,
               bot_ocl_buf,
               wei_ocl_buf,
               top_ocl_buf,
               reserved_ptr, // Unused return_addr.
               R,
               S,
               pad_H, // Like Fwd wino.
               pad_W,
               out_H,
               out_W,
               reserved_ptr, // Unused bias_addr.
               reserved,     // Unused relu_alpha.
               d_buf.byte_stride.nk,
               d_buf.byte_stride.c,
               d_buf.byte_stride.h,
               d_buf.byte_stride.w,
               f_buf.byte_stride.nk,
               f_buf.byte_stride.c,
               f_buf.byte_stride.h,
               f_buf.byte_stride.w,
               o_buf.byte_stride.nk,
               o_buf.byte_stride.c,
               o_buf.byte_stride.h,
               o_buf.byte_stride.w,
               group_cnt,
               d_buf.byte_stride.g,
               f_buf.byte_stride.g,
               o_buf.byte_stride.g);

        elapsed_time = profile_h.GetKernelTime();
    }
#ifdef NDEBUG
    catch(miopen::Exception& ex)
    {
        MIOPEN_LOG_WE(ex.what());
        return -1;
    }
#endif
    return 0;
}

bool ConvBinWinogradRxSf2x3::IsApplicable(const ConvolutionContext& params) const
{
    if(!params.Is2d())
        return false;
    if(!(params.IsFp32() || params.IsFp16()))
        return false;
    if(miopen::IsDisabled(MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3{}))
        return false;
    if(!params.use_asm_kernels)
        return false;
    if(!params.rmv.IsV3())
        return false;

    const auto name = params.GetStream().GetDeviceName();
    if(!(StartsWith(name, "gfx9")))
        return false;
    if(params.IsFp16() && !(StartsWith(name, "gfx906") || StartsWith(name, "gfx908")))
        return false;

    // clang-format off
    if (! ( (params.kernel_stride_w == 1 || params.kernel_stride_w == 2)
        && params.kernel_stride_w == params.kernel_stride_h
        && params.kernel_dilation_w == 1
        && params.kernel_dilation_h == 1
        && params.bias == 0
        && params.in_layout == "NCHW"))
        return false;
    // clang-format on

    const auto n_inputs_per_group  = params.n_inputs / params.group_counts,
               n_outputs_per_group = params.n_outputs / params.group_counts;

    if(params.direction.IsBackwardWrW())
    {
        if(params.kernel_stride_w == 2)
            return false;
        return IsShaderContraintsMet(params.in_height,
                                     params.in_width,
                                     params.kernel_dilation_h,
                                     params.kernel_dilation_w,
                                     params.batch_sz,    // N
                                     n_inputs_per_group, // K
                                     params.out_height,
                                     params.out_width,
                                     params.kernel_size_h,
                                     params.kernel_size_w,
                                     n_outputs_per_group, // C
                                     params);
    }
    else
    {
        return IsShaderContraintsMet(params.kernel_size_h, // RxS
                                     params.kernel_size_w,
                                     params.kernel_stride_h,
                                     params.kernel_stride_w,
                                     n_inputs_per_group,  // C
                                     n_outputs_per_group, // K
                                     params.in_height,    // HxW
                                     params.in_width,
                                     params.out_height, // OHxOW
                                     params.out_width,
                                     params.batch_sz, // N
                                     params);
    }
}

ConvSolution
ConvBinWinogradRxSf2x3::GetSolution(const ConvolutionContext& params,
                                    const PerformanceConfigConvBinWinogradRxSf2x3& config,
                                    const bool disableConfigOverrideFromEnv) const
{
    const auto n_groups = config.n_groups;
    static bool IsWarned;
    if(!IsWarned)
    {
        if(params.GetStream().GetMaxComputeUnits() > MAX_CU_LIMIT)
            MIOPEN_LOG_WE(SolverDbId(*this) << ": GPU has "
                                            << params.GetStream().GetMaxComputeUnits()
                                            << "CUs, but this solver supports max "
                                            << MAX_CU_LIMIT
                                            << "and thus may show sub-optimal performance.");
        IsWarned = true;
    }

    ConvSolution result;

    const PerformanceConfigConvBinWinogradRxSf2x3* pcfg = &config;
    PerformanceConfigConvBinWinogradRxSf2x3 fromEnv;

    if(!disableConfigOverrideFromEnv)
    {
        std::string s;
        const auto p_asciz = miopen::GetStringEnv(MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3_PERF_VALS{});
        if(p_asciz != nullptr)
        {
            s = std::string(p_asciz);
            if(!s.empty()) // else nothing to parse.
            {
                if(!fromEnv.Deserialize(s) || !fromEnv.IsValid(params))
                {
                    MIOPEN_LOG_E("MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3_PERF_VALS: "
                                 "Bad format or invalid for the problem config: "
                                 << s);
                }
                else
                {
                    MIOPEN_LOG_I("Overridden from env: " << fromEnv.ToString());
                    pcfg = &fromEnv;
                }
            }
        }
    }

    KernelInfo kernel;

    kernel.g_wk.push_back(512 * pcfg->GetNGroups() * params.group_counts);
    kernel.g_wk.push_back(1);
    kernel.g_wk.push_back(1);

    kernel.l_wk.push_back(512);
    kernel.l_wk.push_back(1);
    kernel.l_wk.push_back(1);

    KernelBuildParameters options{
        {"ROCM_METADATA_VERSION", 5},
    };
    kernel.comp_options = options.GenerateFor(kbp::GcnAsm{});

    std::string kernel_name    = "miopenSp3AsmConv";
    std::string kernel_file    = "Conv_Winograd";
    std::string kernel_postfix = "_v21_1_0_gfx9";

    if(params.IsFp32())
        kernel_postfix += "_fp32";
    else
    {
        kernel_postfix += "_fp16_dot2_edc";
    }
    if(params.kernel_stride_w == 1)
    {
        kernel_postfix += "_stride1";
    }
    else if(params.kernel_stride_w == 2 && !params.direction.IsBackwardData())
    {
        kernel_postfix += "_stride2";
    }
    else // if(params.kernel_dilation_h == 2)
    {
        kernel_postfix += "_dilation2";
    }

    if(params.group_counts != 1 || params.direction.IsBackwardWrW())
        kernel_postfix += "_group";

    kernel.kernel_name = kernel_name + kernel_postfix;
    kernel.kernel_file = kernel_file + kernel_postfix + ".s";

    result.construction_params.push_back(kernel);

    if(!params.direction.IsBackwardWrW())
    {
        const bool is_forward     = params.direction.IsForward();
        constexpr int F_REVERSE_R = 1 << 0;
        constexpr int F_REVERSE_S = 1 << 1;
        constexpr int F_FLIP_K_C  = 1 << 2;
        // These are not used yet. Nevertheless let's keep as a shader documentation.
        // constexpr int F_FLIP_DATA_N_C = 1 << 3; // Unsupported in f3x2.
        // constexpr int F_FLIP_OUT_N_K = 1 << 4; // Unsupported in f3x2.
        // constexpr int L_F_ADDR_INDIRECT  = 1 << 6;
        // constexpr int L_F_BIAS  = 1 << 7;
        // constexpr int L_F_LEAKY_RELU  = 1 << 8;
        constexpr int L_F_NKC_STRIDES   = 1 << 9;
        constexpr int L_F_GROUP_STRIDES = 1 << 10;
        int reserved                    = 0;
        int* reserved_ptr               = nullptr;
        int ignore;

        int N, C, H, W, K, out_H, out_W, R, S, pad_H, pad_W;
        GetCompiledInParameters(
            params, &N, &C, &H, &W, &K, &ignore, &out_H, &out_W, &R, &S, &pad_H, &pad_W);
        const auto group_cnt = params.group_counts;
        C                    = C / group_cnt;
        K                    = K / group_cnt;
        int flags            = is_forward ? 0 : F_REVERSE_R + F_REVERSE_S + F_FLIP_K_C;
        flags |= L_F_NKC_STRIDES + L_F_GROUP_STRIDES;

        // cppcheck-suppress unreadVariable
        BuffInfo d_buf(GetGroupConvLayout(GetMemLayout_t(params.in_layout), true),
                       N,
                       C,
                       H,
                       W,
                       group_cnt,
                       GetTypeSize(params.in_data_type)),
            // cppcheck-suppress unreadVariable
            o_buf(GetGroupConvLayout(GetMemLayout_t(params.out_layout), true),
                  N,
                  K,
                  out_H,
                  out_W,
                  group_cnt,
                  GetTypeSize(params.out_data_type)),
            // cppcheck-suppress unreadVariable
            f_buf(GetGroupConvLayout(is_forward ? (MemLayout_t::NCHW)
                                                : GetSwappedNCLayout(MemLayout_t::NCHW),
                                     false),
                  K,
                  C,
                  R,
                  S,
                  group_cnt,
                  GetTypeSize(params.weights_data_type));

        result.invoker_factory = [=](std::vector<Kernel> kernels) {
            return [=](const Handle& handle, const boost::any& primitive_params) {
                const auto k         = handle.Run(kernels[0]);
                const auto& data_ctx = boost::any_cast<conv::DataInvokeParams>(primitive_params);
                const auto& tensors  = data_ctx.tensors;

                // clang-format off
                MIOPEN_LOG_I2(" N=" << N << " G=" << group_cnt << " C=" << C << " H=" << H << " W=" << W << " K=" << K
                    << " n_groups=" << n_groups << " flags=" << flags << " R=" << R << " S=" << S
                    << " pad_H=" << pad_H << " pad_W=" << pad_W << " out_H=" << out_H << " out_W=" << out_W
                    << " d_buf.byte_stride.nk=" << d_buf.byte_stride.nk << " d_buf.byte_stride.c=" << d_buf.byte_stride.c
                    << " d_buf.byte_stride.h=" << d_buf.byte_stride.h << " d_buf.byte_stride.w=" << d_buf.byte_stride.w
                    << " f_buf.byte_stride.nk=" << f_buf.byte_stride.nk << " f_buf.byte_stride.c=" << f_buf.byte_stride.c
                    << " f_buf.byte_stride.h=" << f_buf.byte_stride.h << " f_buf.byte_stride.w=" << f_buf.byte_stride.w
                    << " o_buf.byte_stride.nk=" << o_buf.byte_stride.nk << " o_buf.byte_stride.c=" << o_buf.byte_stride.c
                    << " o_buf.byte_stride.h="  << o_buf.byte_stride.h <<  " o_buf.byte_stride.w=" << o_buf.byte_stride.w
                    << " d_buf.byte_stride.g=" << d_buf.byte_stride.g  << " o_buf.byte_stride.g="  << o_buf.byte_stride.g
                    << " f_buf.byte_stride.g=" << f_buf.byte_stride.g); // clang-format on

                k(N,
                  C,
                  H,
                  W,
                  K,
                  n_groups,
                  flags,
                  reserved,
                  tensors.in,
                  tensors.w,
                  tensors.out,
                  reserved_ptr, // Unused return_addr.
                  R,
                  S,
                  pad_H, // Like Fwd wino.
                  pad_W,
                  out_H,
                  out_W,
                  reserved_ptr, // Unused bias_addr.
                  reserved,     // Unused relu_alpha.
                  d_buf.byte_stride.nk,
                  d_buf.byte_stride.c,
                  d_buf.byte_stride.h,
                  d_buf.byte_stride.w,
                  f_buf.byte_stride.nk,
                  f_buf.byte_stride.c,
                  f_buf.byte_stride.h,
                  f_buf.byte_stride.w,
                  o_buf.byte_stride.nk,
                  o_buf.byte_stride.c,
                  o_buf.byte_stride.h,
                  o_buf.byte_stride.w,
                  group_cnt,
                  d_buf.byte_stride.g,
                  f_buf.byte_stride.g,
                  o_buf.byte_stride.g);
            };
        };
    }
    else
    {
        int unused = 0;
        int N, C, H, W, K, out_H, out_W, R, S;
        GetCompiledInParameters(
            params, &C, &K, &R, &S, &N, &unused, &H, &W, &out_H, &out_W, &unused, &unused);
        const auto group_cnt             = params.group_counts;
        static const int F_NKC_STRIDES   = 1 << 9;
        static const int F_GROUP_STRIDES = 1 << 10;
        int flags                        = F_NKC_STRIDES + F_GROUP_STRIDES;
        N                                = N / group_cnt;
        K                                = K / group_cnt;
        int pad_H                        = params.conv_problem.GetConv().GetConvPads()[0];
        int pad_W                        = params.conv_problem.GetConv().GetConvPads()[1];

        BuffInfo d_buf(
            GetGroupConvLayout(GetSwappedNCLayout(GetMemLayout_t(params.in_layout)), true),
            N,
            C,
            H,
            W,
            group_cnt,
            GetTypeSize(params.in_data_type)),
            o_buf(GetGroupConvLayout(GetSwappedNCLayout(GetMemLayout_t(params.out_layout)), false),
                  N,
                  K,
                  out_H,
                  out_W,
                  group_cnt,
                  GetTypeSize(params.out_data_type)),
            f_buf(GetGroupConvLayout(GetSwappedNCLayout(MemLayout_t::NCHW), true),
                  K,
                  C,
                  R,
                  S,
                  group_cnt,
                  GetTypeSize(params.weights_data_type));

        decltype(auto) batch_sz = params.batch_sz;
        decltype(auto) n_inputs = params.n_inputs;

        result.invoker_factory = [=](std::vector<Kernel> kernels) {
            return [=](const Handle& handle, const boost::any& primitive_params) {
                decltype(auto) invoke_params =
                    boost::any_cast<conv::WrWInvokeParams>(primitive_params);
                const auto& tensors = invoke_params.tensors;

                // clang-format off
                MIOPEN_LOG_I2(" N=" << N << " G=" << group_cnt << " C=" << C << " H=" << H << " W=" << W << " K=" << K
                    << " n_groups=" << n_groups << " flags=" << flags << " R=" << R << " S=" << S
                    << " pad_H=" << pad_H << " pad_W=" << pad_W << " out_H=" << out_H << " out_W=" << out_W
                    << " d_buf.byte_stride.nk=" << d_buf.byte_stride.nk << " d_buf.byte_stride.c=" << d_buf.byte_stride.c
                    << " d_buf.byte_stride.h=" << d_buf.byte_stride.h << " d_buf.byte_stride.w=" << d_buf.byte_stride.w
                    << " f_buf.byte_stride.nk=" << f_buf.byte_stride.nk << " f_buf.byte_stride.c=" << f_buf.byte_stride.c
                    << " f_buf.byte_stride.h=" << f_buf.byte_stride.h << " f_buf.byte_stride.w=" << f_buf.byte_stride.w
                    << " o_buf.byte_stride.nk=" << o_buf.byte_stride.nk << " o_buf.byte_stride.c=" << o_buf.byte_stride.c
                    << " o_buf.byte_stride.h="  << o_buf.byte_stride.h <<  " o_buf.byte_stride.w=" << o_buf.byte_stride.w
                    << " d_buf.byte_stride.g=" << d_buf.byte_stride.g  << " o_buf.byte_stride.g="  << o_buf.byte_stride.g
                    << " f_buf.byte_stride.g=" << f_buf.byte_stride.g); // clang-format on
                MIOPEN_LOG_I2(" ctx.batch_sz=" << batch_sz << "ctx.n_inputs=" << n_inputs);

                int reserved      = 0;
                int* reserved_ptr = nullptr;

                handle.Run(kernels[0])(N,
                                       C,
                                       H,
                                       W,
                                       K,
                                       n_groups,
                                       flags,
                                       reserved,
                                       tensors.x,
                                       tensors.dy,
                                       tensors.dw,
                                       reserved_ptr, // Unused return_addr.
                                       R,
                                       S,
                                       pad_H, // Like Fwd wino.
                                       pad_W,
                                       out_H,
                                       out_W,
                                       reserved_ptr, // Unused bias_addr.
                                       reserved,     // Unused relu_alpha.
                                       d_buf.byte_stride.nk,
                                       d_buf.byte_stride.c,
                                       d_buf.byte_stride.h,
                                       d_buf.byte_stride.w,
                                       f_buf.byte_stride.nk,
                                       f_buf.byte_stride.c,
                                       f_buf.byte_stride.h,
                                       f_buf.byte_stride.w,
                                       o_buf.byte_stride.nk,
                                       o_buf.byte_stride.c,
                                       o_buf.byte_stride.h,
                                       o_buf.byte_stride.w,
                                       group_cnt,
                                       d_buf.byte_stride.g,
                                       f_buf.byte_stride.g,
                                       o_buf.byte_stride.g);
            };
        };
    }

    return result;
}

} // namespace solver
} // namespace miopen

max_stars_repo_name: ashishfarmer/MIOpen

max_issues_repo_name: ashishfarmer/MIOpen

max_forks_repo_name: ashishfarmer/MIOpen


text:
#ifndef UUID_98C06485_8AE9_47DA_B99F_62CA5AF00FF4
#define UUID_98C06485_8AE9_47DA_B99F_62CA5AF00FF4
#pragma once

#include "utils.hpp"
#include <atomic>
#include <iostream>
#include <cstdint>
#include <memory>
#include <typeindex>
#include <vector>
#include <boost/intrusive_ptr.hpp>

namespace Neptools
{

template <typename Item, typename... Rest> struct IndexOf;
template <typename Item, typename... Rest>
struct IndexOf<Item, Item, Rest...>
    : std::integral_constant<size_t, 0> {};

template <typename Item, typename Head, typename... Rest>
struct IndexOf<Item, Head, Rest...>
    : std::integral_constant<size_t, 1 + IndexOf<Item, Rest...>::value> {};

template <typename Item, typename... Args>
constexpr auto IndexOfV = IndexOf<Item, Args...>::value;

NEPTOOLS_STATIC_ASSERT(IndexOfV<int, float, double, int> == 2);

template <typename... Args>
class DynamicStruct
{
    template <typename Ret, size_t I, typename T, typename... TRest,
              typename Thiz, typename Fun, typename... FunArgs>
    static Ret VisitHlp(Thiz thiz, size_t i, Fun&& fun, FunArgs&&... args)
    {
        if (thiz->type->items[i].idx == I)
            return fun(thiz->template Get<T>(i), thiz->GetSize(i),
                       std::forward<FunArgs>(args)...);
        else
            return VisitHlp<Ret, I+1, TRest...>(
                thiz, i, std::forward<Fun>(fun), std::forward<FunArgs>(args)...);
    }

    template <typename Ret, size_t I, typename Thiz, typename Fun,
              typename... FunArgs>
    static Ret VisitHlp(Thiz, size_t, Fun&&, FunArgs&&...) { abort(); }

public:
    template <typename T>
    static constexpr size_t GetIndexFromType()
    { return IndexOfV<T, Args...>; }

    static constexpr const size_t SIZE_OF[] = { sizeof(Args)... };
    static constexpr const size_t ALIGN_OF[] = { alignof(Args)... };

    struct Type
    {
        Type(const Type&) = delete;
        void operator=(const Type&) = delete;
        ~Type() = delete;

        mutable std::atomic<size_t> refcount;
        size_t item_count, size;
        struct Item
        {
            size_t idx;
            size_t size;
            size_t offset;
        };
        Item items[1];
    };
    using TypePtr = boost::intrusive_ptr<const Type>;

    friend void intrusive_ptr_add_ref(const Type* t)
    { t->refcount.fetch_add(1, std::memory_order_relaxed); }

    friend void intrusive_ptr_release(const Type* t)
    {
        if (t->refcount.fetch_sub(1, std::memory_order_acq_rel) == 1)
            ::operator delete(const_cast<Type*>(t));
    }

    class TypeBuilder
    {
    public:
        auto& GetDesc() { return desc; }
        const auto& GetDesc() const { return desc; }

        template <typename T>
        void Add(size_t size = sizeof(T))
        { desc.emplace_back(IndexOfV<T, Args...>, size); }

        void Add(size_t i, size_t size)
        {
            NEPTOOLS_ASSERT_MSG(i < sizeof...(Args), "index out of range");
            desc.push_back({i, size});
        }

        TypePtr Build() const
        {
            auto ptr = operator new(
                sizeof(Type) + (desc.size() - 1) * sizeof(typename Type::Item));
            boost::intrusive_ptr<Type> ret{static_cast<Type*>(ptr), false};
            ret->refcount.store(1, std::memory_order_relaxed);
            ret->item_count = desc.size();

            size_t offs = 0;
            for (size_t i = 0; i < desc.size(); ++i)
            {
                ret->items[i].idx = desc[i].first;
                ret->items[i].size = desc[i].second;
                ret->items[i].offset = offs;

                offs += desc[i].second;
                auto al = ALIGN_OF[desc[i].first];
                offs = (offs + al - 1) / al * al;
            }
            ret->size = offs;

            return ret;
        }

    private:
        std::vector<std::pair<size_t, size_t>> desc;
    };

    // actual class begin
    DynamicStruct() = default;
    DynamicStruct(const TypePtr& type) : type{type}
    {
        data.reset(new char[type->size]);
        ForEach([](auto& x, size_t)
                { new (&x) std::remove_reference_t<decltype(x)>; });
    }

    DynamicStruct(DynamicStruct&& o) : type{o.type}, data{std::move(o.data)}
    { o.type = nullptr; o.data = nullptr; }
    DynamicStruct& operator=(DynamicStruct&& o)
    {
        if (this != &o)
        {
            type = o.type;
            data = std::move(o.data);
            o.type = nullptr;
            o.data = nullptr;
        }
        return *this;
    }
    ~DynamicStruct()
    {
        ForEach([](auto& x, size_t)
                {
                    using T = std::remove_reference_t<decltype(x)>;
                    x.~T();
                });
    }

    size_t GetSize() const { return type->item_count; }
    size_t GetSize(size_t i) const
    {
        NEPTOOLS_ASSERT_MSG(i < GetSize(), "index out of range");
        return type->items[i].size;
    }
    size_t GetTypeIndex(size_t i) const
    {
        NEPTOOLS_ASSERT_MSG(i < GetSize(), "index out of range");
        return type->items[i].idx;
    }

    template <typename T>
    bool Is(size_t i) const
    {
        return GetTypeIndex(i) == GetIndexFromType<T>();
    }

    TypePtr GetRawType() const { return type; }
    void* GetData() { return data.get(); }
    const void* GetData() const { return data.get(); }

    template <typename T>
    T& Get(size_t i)
    {
        NEPTOOLS_ASSERT_MSG(Is<T>(i), "specified item is not T");
        return *reinterpret_cast<T*>(&data[type->items[i].offset]);
    }

    template <typename T>
    const T& Get(size_t i) const
    {
        NEPTOOLS_ASSERT_MSG(Is<T>(i), "specified item is not T");
        return *reinterpret_cast<const T*>(&data[type->items[i].offset]);
    }

    template <typename Ret = void, typename... FunArgs>
    Ret Visit(size_t i, FunArgs&&... f)
    { return VisitHlp<Ret, 0, Args...>(this, i, std::forward<FunArgs>(f)...); }

    template <typename... FunArgs>
    void ForEach(FunArgs&&... f)
    {
        for (size_t i = 0; i < type->item_count; ++i)
            Visit(i, std::forward<FunArgs>(f)...);
    }

    // const version
    template <typename Ret = void, typename... FunArgs>
    Ret Visit(size_t i, FunArgs&&... f) const
    { return VisitHlp<Ret, 0, Args...>(this, i, std::forward<FunArgs>(f)...); }

    template <typename... FunArgs>
    void ForEach(FunArgs&&... args) const
    {
        for (size_t i = 0; i < type->item_count; ++i)
            Visit(i, std::forward<FunArgs>(args)...);
    }

private:
    TypePtr type;
    std::unique_ptr<char[]> data;
};

template <typename... Args>
constexpr const size_t DynamicStruct<Args...>::SIZE_OF[];

template <typename... Args>
constexpr const size_t DynamicStruct<Args...>::ALIGN_OF[];

}
#endif

max_stars_repo_name: clienthax/neptools

max_issues_repo_name: clienthax/neptools

max_forks_repo_name: clienthax/neptools