From 1f17c569b961123ef97def7bfebc1266a2e4cd79 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pawe=C5=82=20Bylica?= Date: Wed, 27 May 2015 13:23:00 +0200 Subject: [PATCH 01/20] Change VM interface to return a copy of output. --- evmjit/libevmjit-cpp/JitVM.cpp | 4 ++-- evmjit/libevmjit-cpp/JitVM.h | 2 +- libethereum/Executive.cpp | 12 +++++------- libethereum/Executive.h | 5 ++--- libevm/SmartVM.cpp | 4 ++-- libevm/SmartVM.h | 2 +- libevm/VM.cpp | 2 +- libevm/VM.h | 2 +- libevm/VMFace.h | 7 ++++++- test/fuzzTesting/checkRandomVMTest.cpp | 4 ++-- test/fuzzTesting/createRandomVMTest.cpp | 2 +- test/libevm/vm.cpp | 4 +--- 12 files changed, 25 insertions(+), 25 deletions(-) diff --git a/evmjit/libevmjit-cpp/JitVM.cpp b/evmjit/libevmjit-cpp/JitVM.cpp index 5193168a4..ac8df545f 100644 --- a/evmjit/libevmjit-cpp/JitVM.cpp +++ b/evmjit/libevmjit-cpp/JitVM.cpp @@ -18,7 +18,7 @@ namespace eth extern "C" void env_sload(); // fake declaration for linker symbol stripping workaround, see a call below -bytesConstRef JitVM::go(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp, uint64_t _step) +bytesConstRef JitVM::execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp, uint64_t _step) { using namespace jit; @@ -33,7 +33,7 @@ bytesConstRef JitVM::go(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp, ui { cwarn << "Execution rejected by EVM JIT (gas limit: " << io_gas << "), executing with interpreter"; m_fallbackVM = VMFactory::create(VMKind::Interpreter); - return m_fallbackVM->go(io_gas, _ext, _onOp, _step); + return m_fallbackVM->execImpl(io_gas, _ext, _onOp, _step); } m_data.gas = static_cast(io_gas); diff --git a/evmjit/libevmjit-cpp/JitVM.h b/evmjit/libevmjit-cpp/JitVM.h index e6864f885..7c8b4ceb8 100644 --- a/evmjit/libevmjit-cpp/JitVM.h +++ b/evmjit/libevmjit-cpp/JitVM.h @@ -11,7 +11,7 @@ namespace eth class JitVM: public VMFace { public: - virtual bytesConstRef go(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp = {}, uint64_t _steps = (uint64_t)-1) override final; + virtual bytesConstRef execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp = {}, uint64_t _steps = (uint64_t)-1) override final; private: jit::RuntimeData m_data; diff --git a/libethereum/Executive.cpp b/libethereum/Executive.cpp index d7bad0ff9..122432d85 100644 --- a/libethereum/Executive.cpp +++ b/libethereum/Executive.cpp @@ -46,7 +46,7 @@ u256 Executive::gasUsed() const ExecutionResult Executive::executionResult() const { - return ExecutionResult(gasUsed(), m_excepted, m_newAddress, m_out, m_codeDeposit, m_ext ? m_ext->sub.refunds : 0, m_depositSize, m_gasForDeposit); + return ExecutionResult(gasUsed(), m_excepted, m_newAddress, out(), m_codeDeposit, m_ext ? m_ext->sub.refunds : 0, m_depositSize, m_gasForDeposit); } void Executive::accrueSubState(SubState& _parentContext) @@ -145,8 +145,7 @@ bool Executive::call(CallParameters const& _p, u256 const& _gasPrice, Address co else { m_gas = (u256)(_p.gas - g); - m_precompiledOut = it->second.exec(_p.data); - m_out = &m_precompiledOut; + m_out = it->second.exec(_p.data); } } else @@ -219,7 +218,7 @@ bool Executive::go(OnOpFunc const& _onOp) #endif try { - m_out = m_vm->go(m_gas, *m_ext, _onOp); + m_out = m_vm->exec(m_gas, *m_ext, _onOp); if (m_isCreation) { @@ -232,11 +231,10 @@ bool Executive::go(OnOpFunc const& _onOp) } else { - m_codeDeposit = CodeDeposit::Failed; - m_out.reset(); + m_out.clear(); } - m_s.m_cache[m_newAddress].setCode(m_out.toBytes()); + m_s.m_cache[m_newAddress].setCode(std::move(m_out)); } } catch (StepsDone const&) diff --git a/libethereum/Executive.h b/libethereum/Executive.h index 04d5857b3..bd9fb178e 100644 --- a/libethereum/Executive.h +++ b/libethereum/Executive.h @@ -109,7 +109,7 @@ public: /// @returns gas remaining after the transaction/operation. Valid after the transaction has been executed. u256 gas() const { return m_gas; } /// @returns output data of the transaction/operation. - bytesConstRef out() const { return m_out; } + bytesConstRef out() const { return {m_out.data(), m_out.size()}; } /// @returns the new address for the created contract in the CREATE operation. h160 newAddress() const { return m_newAddress; } /// @returns true iff the operation ended with a VM exception. @@ -123,8 +123,7 @@ private: LastHashes m_lastHashes; std::shared_ptr m_ext; ///< The VM externality object for the VM execution or null if no VM is required. std::unique_ptr m_vm; ///< The VM object or null if no VM is required. - bytes m_precompiledOut; ///< Used for the output when there is no VM for a contract (i.e. precompiled). - bytesConstRef m_out; ///< The copyable output. + bytes m_out; ///< The VM execution output. Address m_newAddress; ///< The address of the created contract in the case of create() being called. unsigned m_depth = 0; ///< The context's call-depth. diff --git a/libevm/SmartVM.cpp b/libevm/SmartVM.cpp index d37abc1bf..27263fc2d 100644 --- a/libevm/SmartVM.cpp +++ b/libevm/SmartVM.cpp @@ -41,7 +41,7 @@ namespace } } -bytesConstRef SmartVM::go(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp, uint64_t _steps) +bytesConstRef SmartVM::execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp, uint64_t _steps) { auto codeHash = sha3(_ext.code); auto vmKind = VMKind::Interpreter; // default VM @@ -68,7 +68,7 @@ bytesConstRef SmartVM::go(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp, // TODO: Selected VM must be kept only because it returns reference to its internal memory. // VM implementations should be stateless, without escaping memory reference. m_selectedVM = VMFactory::create(vmKind); - return m_selectedVM->go(io_gas, _ext, _onOp, _steps); + return m_selectedVM->execImpl(io_gas, _ext, _onOp, _steps); } } diff --git a/libevm/SmartVM.h b/libevm/SmartVM.h index cc198c0c2..01e163cd2 100644 --- a/libevm/SmartVM.h +++ b/libevm/SmartVM.h @@ -31,7 +31,7 @@ namespace eth class SmartVM: public VMFace { public: - virtual bytesConstRef go(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp = {}, uint64_t _steps = (uint64_t)-1) override final; + virtual bytesConstRef execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp = {}, uint64_t _steps = (uint64_t)-1) override final; private: std::unique_ptr m_selectedVM; diff --git a/libevm/VM.cpp b/libevm/VM.cpp index df8432ac6..3a4124d67 100644 --- a/libevm/VM.cpp +++ b/libevm/VM.cpp @@ -45,7 +45,7 @@ static array metrics() return s_ret; } -bytesConstRef VM::go(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp, uint64_t _steps) +bytesConstRef VM::execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp, uint64_t _steps) { // Reset leftovers from possible previous run m_curPC = 0; diff --git a/libevm/VM.h b/libevm/VM.h index 99f608227..0a33e9fae 100644 --- a/libevm/VM.h +++ b/libevm/VM.h @@ -52,7 +52,7 @@ inline u256 fromAddress(Address _a) class VM: public VMFace { public: - virtual bytesConstRef go(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp = {}, uint64_t _steps = (uint64_t)-1) override final; + virtual bytesConstRef execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp = {}, uint64_t _steps = (uint64_t)-1) override final; u256 curPC() const { return m_curPC; } diff --git a/libevm/VMFace.h b/libevm/VMFace.h index d2a12e0ca..8fd7ebcb3 100644 --- a/libevm/VMFace.h +++ b/libevm/VMFace.h @@ -43,7 +43,12 @@ public: VMFace(VMFace const&) = delete; VMFace& operator=(VMFace const&) = delete; - virtual bytesConstRef go(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp = {}, uint64_t _steps = (uint64_t)-1) = 0; + bytes exec(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp = {}, uint64_t _steps = (uint64_t)-1) + { + return execImpl(io_gas, _ext, _onOp, _steps).toVector(); + } + + virtual bytesConstRef execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp = {}, uint64_t _steps = (uint64_t)-1) = 0; }; } diff --git a/test/fuzzTesting/checkRandomVMTest.cpp b/test/fuzzTesting/checkRandomVMTest.cpp index 13d677e17..a6ade07f1 100644 --- a/test/fuzzTesting/checkRandomVMTest.cpp +++ b/test/fuzzTesting/checkRandomVMTest.cpp @@ -98,9 +98,9 @@ bool doVMTest(mValue& _v) try { auto vm = eth::VMFactory::create(); - output = vm->go(fev.gas, fev, fev.simpleTrace()).toBytes(); + output = vm->exec(fev.gas, fev, fev.simpleTrace()); } - catch (eth::VMException) + catch (eth::VMException const&) { cnote << "Safe VM Exception"; vmExceptionOccured = true; diff --git a/test/fuzzTesting/createRandomVMTest.cpp b/test/fuzzTesting/createRandomVMTest.cpp index 3196590d4..93040d6de 100644 --- a/test/fuzzTesting/createRandomVMTest.cpp +++ b/test/fuzzTesting/createRandomVMTest.cpp @@ -160,7 +160,7 @@ void doMyTests(json_spirit::mValue& _v) bool vmExceptionOccured = false; try { - output = vm->go(fev.gas, fev, fev.simpleTrace()).toBytes(); + output = vm->exec(fev.gas, fev, fev.simpleTrace()); } catch (eth::VMException const& _e) { diff --git a/test/libevm/vm.cpp b/test/libevm/vm.cpp index cb63e993c..c407c8184 100644 --- a/test/libevm/vm.cpp +++ b/test/libevm/vm.cpp @@ -329,12 +329,10 @@ void doVMTests(json_spirit::mValue& v, bool _fillin) { auto vm = eth::VMFactory::create(); auto vmtrace = Options::get().vmtrace ? fev.simpleTrace() : OnOpFunc{}; - auto outputRef = bytesConstRef{}; { Listener::ExecTimeGuard guard{i.first}; - outputRef = vm->go(fev.gas, fev, vmtrace); + output = vm->exec(fev.gas, fev, vmtrace); } - output = outputRef.toBytes(); } catch (VMException const&) { From 72856a8af658852206f6e117a71b4eb0f6375a44 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pawe=C5=82=20Bylica?= Date: Thu, 28 May 2015 08:56:21 +0200 Subject: [PATCH 02/20] Change the way execution results are collected. Changes handling ExecutionResult by Executive. From now execution results are collected on if a storage for results (ExecutionResult) is provided to an Executiove instance up front. This change allow better output management for calls - VM interface improved. --- evmjit/libevmjit-cpp/JitVM.h | 2 +- libethereum/Client.cpp | 3 +- libethereum/Executive.cpp | 60 ++++++++++++------- libethereum/Executive.h | 14 ++--- libethereum/ExtVM.cpp | 1 - libethereum/State.cpp | 4 +- libethereum/Transaction.h | 17 +----- libevm/SmartVM.h | 2 +- libevm/VM.h | 2 +- libevm/VMFace.h | 12 +++- mix/MixClient.cpp | 5 +- test/libsolidity/solidityExecutionFramework.h | 4 +- 12 files changed, 72 insertions(+), 54 deletions(-) diff --git a/evmjit/libevmjit-cpp/JitVM.h b/evmjit/libevmjit-cpp/JitVM.h index 7c8b4ceb8..03c3c8880 100644 --- a/evmjit/libevmjit-cpp/JitVM.h +++ b/evmjit/libevmjit-cpp/JitVM.h @@ -11,7 +11,7 @@ namespace eth class JitVM: public VMFace { public: - virtual bytesConstRef execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp = {}, uint64_t _steps = (uint64_t)-1) override final; + virtual bytesConstRef execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp, uint64_t _steps) override final; private: jit::RuntimeData m_data; diff --git a/libethereum/Client.cpp b/libethereum/Client.cpp index 6dec81b38..d9813c725 100644 --- a/libethereum/Client.cpp +++ b/libethereum/Client.cpp @@ -440,9 +440,10 @@ ExecutionResult Client::call(Address _dest, bytes const& _data, u256 _gas, u256 temp = m_postMine; temp.addBalance(_from, _value + _gasPrice * _gas); Executive e(temp, LastHashes(), 0); + e.setResultRef(ret); if (!e.call(_dest, _from, _value, _gasPrice, &_data, _gas)) e.go(); - ret = e.executionResult(); + e.finalize(); } catch (...) { diff --git a/libethereum/Executive.cpp b/libethereum/Executive.cpp index 122432d85..fc8ba22a3 100644 --- a/libethereum/Executive.cpp +++ b/libethereum/Executive.cpp @@ -44,11 +44,6 @@ u256 Executive::gasUsed() const return m_t.gas() - m_gas; } -ExecutionResult Executive::executionResult() const -{ - return ExecutionResult(gasUsed(), m_excepted, m_newAddress, out(), m_codeDeposit, m_ext ? m_ext->sub.refunds : 0, m_depositSize, m_gasForDeposit); -} - void Executive::accrueSubState(SubState& _parentContext) { if (m_ext) @@ -145,7 +140,8 @@ bool Executive::call(CallParameters const& _p, u256 const& _gasPrice, Address co else { m_gas = (u256)(_p.gas - g); - m_out = it->second.exec(_p.data); + auto out = it->second.exec(_p.data); // FIXME: Avoid double copy + bytesConstRef{&out}.copyTo(_p.out); } } else @@ -153,7 +149,7 @@ bool Executive::call(CallParameters const& _p, u256 const& _gasPrice, Address co m_gas = _p.gas; if (m_s.addressHasCode(_p.codeAddress)) { - m_vm = VMFactory::create(); + m_outRef = _p.out; // Save ref to expected output buffer to be used in go() bytes const& c = m_s.code(_p.codeAddress); m_ext = make_shared(m_s, m_lastHashes, _p.receiveAddress, _p.senderAddress, _origin, _p.value, _gasPrice, _p.data, &c, m_depth); } @@ -175,10 +171,7 @@ bool Executive::create(Address _sender, u256 _endowment, u256 _gasPrice, u256 _g // Execute _init. if (!_init.empty()) - { - m_vm = VMFactory::create(); m_ext = make_shared(m_s, m_lastHashes, m_newAddress, _sender, _origin, _endowment, _gasPrice, bytesConstRef(), _init, m_depth); - } m_s.m_cache[m_newAddress] = Account(m_s.balance(m_newAddress), Account::ContractConception); m_s.transferBalance(_sender, m_newAddress, _endowment); @@ -211,30 +204,47 @@ OnOpFunc Executive::simpleTrace() bool Executive::go(OnOpFunc const& _onOp) { - if (m_vm) + if (m_ext) { #if ETH_TIMED_EXECUTIONS boost::timer t; #endif try { - m_out = m_vm->exec(m_gas, *m_ext, _onOp); - + auto vm = VMFactory::create(); if (m_isCreation) { - m_gasForDeposit = m_gas; - m_depositSize = m_out.size(); - if (m_out.size() * c_createDataGas <= m_gas) + auto out = vm->exec(m_gas, *m_ext, _onOp); + if (m_res) + { + m_res->gasForDeposit = m_gas; + m_res->depositSize = out.size(); + } + if (out.size() * c_createDataGas <= m_gas) { - m_codeDeposit = CodeDeposit::Success; - m_gas -= m_out.size() * c_createDataGas; + if (m_res) + m_res->codeDeposit = CodeDeposit::Success; + m_gas -= out.size() * c_createDataGas; } else { - m_codeDeposit = CodeDeposit::Failed; - m_out.clear(); + if (m_res) + m_res->codeDeposit = CodeDeposit::Failed; + out.clear(); } - m_s.m_cache[m_newAddress].setCode(std::move(m_out)); + if (m_res) + m_res->output = out; // copy output to execution result + m_s.m_cache[m_newAddress].setCode(std::move(out)); // FIXME: Set only if Success? + } + else + { + if (m_res) + { + m_res->output = vm->exec(m_gas, *m_ext, _onOp); // take full output + bytesConstRef{&m_res->output}.copyTo(m_outRef); + } + else + vm->exec(m_gas, *m_ext, m_outRef, _onOp); // take only expected output } } catch (StepsDone const&) @@ -293,4 +303,12 @@ void Executive::finalize() // Logs.. if (m_ext) m_logs = m_ext->sub.logs; + + if (m_res) // Collect results + { + m_res->gasUsed = gasUsed(); + m_res->excepted = m_excepted; // TODO: m_except is used only in ExtVM::call + m_res->newAddress = m_newAddress; + m_res->gasRefunded = m_ext ? m_ext->sub.refunds : 0; + } } diff --git a/libethereum/Executive.h b/libethereum/Executive.h index bd9fb178e..8b2906b9b 100644 --- a/libethereum/Executive.h +++ b/libethereum/Executive.h @@ -108,29 +108,25 @@ public: /// @returns gas remaining after the transaction/operation. Valid after the transaction has been executed. u256 gas() const { return m_gas; } - /// @returns output data of the transaction/operation. - bytesConstRef out() const { return {m_out.data(), m_out.size()}; } + /// @returns the new address for the created contract in the CREATE operation. h160 newAddress() const { return m_newAddress; } /// @returns true iff the operation ended with a VM exception. bool excepted() const { return m_excepted != TransactionException::None; } - /// Get the above in an amalgamated fashion. - ExecutionResult executionResult() const; + /// Collect execution results in the result storage provided. + void setResultRef(ExecutionResult& _res) { m_res = &_res; } private: State& m_s; ///< The state to which this operation/transaction is applied. LastHashes m_lastHashes; std::shared_ptr m_ext; ///< The VM externality object for the VM execution or null if no VM is required. - std::unique_ptr m_vm; ///< The VM object or null if no VM is required. - bytes m_out; ///< The VM execution output. + bytesRef m_outRef; ///< Reference to "expected output" buffer. + ExecutionResult* m_res = nullptr; ///< Optional storage for execution results. Address m_newAddress; ///< The address of the created contract in the case of create() being called. unsigned m_depth = 0; ///< The context's call-depth. bool m_isCreation = false; ///< True if the transaction creates a contract, or if create() is called. - unsigned m_depositSize = 0; ///< Amount of code of the creation's attempted deposit. - u256 m_gasForDeposit; ///< Amount of gas remaining for the code deposit phase. - CodeDeposit m_codeDeposit = CodeDeposit::None; ///< True if an attempted deposit failed due to lack of gas. TransactionException m_excepted = TransactionException::None; ///< Details if the VM's execution resulted in an exception. u256 m_gas = 0; ///< The gas for EVM code execution. Initial amount before go() execution, final amount after go() execution. diff --git a/libethereum/ExtVM.cpp b/libethereum/ExtVM.cpp index 488d8e4b7..782d4d79c 100644 --- a/libethereum/ExtVM.cpp +++ b/libethereum/ExtVM.cpp @@ -35,7 +35,6 @@ bool ExtVM::call(CallParameters& _p) e.accrueSubState(sub); } _p.gas = e.gas(); - e.out().copyTo(_p.out); return !e.excepted(); } diff --git a/libethereum/State.cpp b/libethereum/State.cpp index c753f57ea..277e9a649 100644 --- a/libethereum/State.cpp +++ b/libethereum/State.cpp @@ -1125,6 +1125,8 @@ ExecutionResult State::execute(LastHashes const& _lh, Transaction const& _t, Per // Create and initialize the executive. This will throw fairly cheaply and quickly if the // transaction is bad in any way. Executive e(*this, _lh, 0); + ExecutionResult res; + e.setResultRef(res); e.initialize(_t); // Uncommitting is a non-trivial operation - only do it once we've verified as much of the @@ -1181,7 +1183,7 @@ ExecutionResult State::execute(LastHashes const& _lh, Transaction const& _t, Per m_transactionSet.insert(e.t().sha3()); } - return e.executionResult(); + return res; } State State::fromPending(unsigned _i) const diff --git a/libethereum/Transaction.h b/libethereum/Transaction.h index 09d6cd54c..3a78ced26 100644 --- a/libethereum/Transaction.h +++ b/libethereum/Transaction.h @@ -74,25 +74,14 @@ TransactionException toTransactionException(VMException const& _e); /// Description of the result of executing a transaction. struct ExecutionResult { - ExecutionResult() = default; - ExecutionResult(u256 const& _gasUsed, TransactionException _excepted, Address const& _newAddress, bytesConstRef _output, CodeDeposit _codeDeposit, u256 const& _gasRefund, unsigned _depositSize, u256 const& _gasForDeposit): - gasUsed(_gasUsed), - excepted(_excepted), - newAddress(_newAddress), - output(_output.toBytes()), - codeDeposit(_codeDeposit), - gasRefunded(_gasRefund), - depositSize(_depositSize), - gasForDeposit(_gasForDeposit) - {} u256 gasUsed = 0; TransactionException excepted = TransactionException::Unknown; Address newAddress; bytes output; - CodeDeposit codeDeposit = CodeDeposit::None; + CodeDeposit codeDeposit = CodeDeposit::None; ///< Failed if an attempted deposit failed due to lack of gas. u256 gasRefunded = 0; - unsigned depositSize = 0; - u256 gasForDeposit; + unsigned depositSize = 0; ///< Amount of code of the creation's attempted deposit. + u256 gasForDeposit; ///< Amount of gas remaining for the code deposit phase. }; std::ostream& operator<<(std::ostream& _out, ExecutionResult const& _er); diff --git a/libevm/SmartVM.h b/libevm/SmartVM.h index 01e163cd2..73d515f29 100644 --- a/libevm/SmartVM.h +++ b/libevm/SmartVM.h @@ -31,7 +31,7 @@ namespace eth class SmartVM: public VMFace { public: - virtual bytesConstRef execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp = {}, uint64_t _steps = (uint64_t)-1) override final; + virtual bytesConstRef execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp, uint64_t _steps) override final; private: std::unique_ptr m_selectedVM; diff --git a/libevm/VM.h b/libevm/VM.h index 0a33e9fae..da3a843ed 100644 --- a/libevm/VM.h +++ b/libevm/VM.h @@ -52,7 +52,7 @@ inline u256 fromAddress(Address _a) class VM: public VMFace { public: - virtual bytesConstRef execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp = {}, uint64_t _steps = (uint64_t)-1) override final; + virtual bytesConstRef execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp, uint64_t _steps) override final; u256 curPC() const { return m_curPC; } diff --git a/libevm/VMFace.h b/libevm/VMFace.h index 8fd7ebcb3..7a344a072 100644 --- a/libevm/VMFace.h +++ b/libevm/VMFace.h @@ -43,12 +43,22 @@ public: VMFace(VMFace const&) = delete; VMFace& operator=(VMFace const&) = delete; + /// Execute EVM code by VM. + /// + /// @param _out Expected output + void exec(u256& io_gas, ExtVMFace& _ext, bytesRef _out, OnOpFunc const& _onOp = {}, uint64_t _steps = (uint64_t)-1) + { + execImpl(io_gas, _ext, _onOp, _steps).copyTo(_out); + } + + /// The same as above but returns a copy of full output. bytes exec(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp = {}, uint64_t _steps = (uint64_t)-1) { return execImpl(io_gas, _ext, _onOp, _steps).toVector(); } - virtual bytesConstRef execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp = {}, uint64_t _steps = (uint64_t)-1) = 0; + /// VM implementation + virtual bytesConstRef execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp, uint64_t _steps) = 0; }; } diff --git a/mix/MixClient.cpp b/mix/MixClient.cpp index b011e0cec..2644a8c65 100644 --- a/mix/MixClient.cpp +++ b/mix/MixClient.cpp @@ -130,7 +130,9 @@ void MixClient::executeTransaction(Transaction const& _t, State& _state, bool _c State execState = _state; execState.addBalance(t.sender(), t.gas() * t.gasPrice()); //give it enough balance for gas estimation + eth::ExecutionResult er; Executive execution(execState, lastHashes, 0); + execution.setResultRef(er); execution.initialize(t); execution.execute(); std::vector machineStates; @@ -186,7 +188,6 @@ void MixClient::executeTransaction(Transaction const& _t, State& _state, bool _c execution.go(onOp); execution.finalize(); - dev::eth::ExecutionResult er = execution.executionResult(); switch (er.excepted) { @@ -213,7 +214,7 @@ void MixClient::executeTransaction(Transaction const& _t, State& _state, bool _c }; ExecutionResult d; - d.result = execution.executionResult(); + d.result = er; d.machineStates = machineStates; d.executionCode = std::move(codes); d.transactionData = std::move(data); diff --git a/test/libsolidity/solidityExecutionFramework.h b/test/libsolidity/solidityExecutionFramework.h index 0f80e9f59..ea277421a 100644 --- a/test/libsolidity/solidityExecutionFramework.h +++ b/test/libsolidity/solidityExecutionFramework.h @@ -135,6 +135,8 @@ protected: { m_state.addBalance(m_sender, _value); // just in case eth::Executive executive(m_state, eth::LastHashes(), 0); + eth::ExecutionResult res; + executive.setResultRef(res); eth::Transaction t = _isCreation ? eth::Transaction(_value, m_gasPrice, m_gas, _data, 0, KeyPair::create().sec()) : eth::Transaction(_value, m_gasPrice, m_gas, m_contractAddress, _data, 0, KeyPair::create().sec()); bytes transactionRLP = t.rlp(); @@ -161,7 +163,7 @@ protected: m_state.noteSending(m_sender); executive.finalize(); m_gasUsed = executive.gasUsed(); - m_output = executive.out().toVector(); + m_output = std::move(res.output); // FIXME: Looks like Framework needs ExecutiveResult embedded m_logs = executive.logs(); } From 9891028cb8bed067559c2b44d9d4a1a0180676f0 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pawe=C5=82=20Bylica?= Date: Thu, 28 May 2015 09:49:48 +0200 Subject: [PATCH 03/20] Improve output handling in precompiled contracts. --- libethereum/Executive.cpp | 3 +- libethereum/Precompiled.cpp | 56 ++++++++++++++++++++----------------- libethereum/Precompiled.h | 2 +- 3 files changed, 32 insertions(+), 29 deletions(-) diff --git a/libethereum/Executive.cpp b/libethereum/Executive.cpp index fc8ba22a3..e33195631 100644 --- a/libethereum/Executive.cpp +++ b/libethereum/Executive.cpp @@ -140,8 +140,7 @@ bool Executive::call(CallParameters const& _p, u256 const& _gasPrice, Address co else { m_gas = (u256)(_p.gas - g); - auto out = it->second.exec(_p.data); // FIXME: Avoid double copy - bytesConstRef{&out}.copyTo(_p.out); + it->second.exec(_p.data, _p.out); } } else diff --git a/libethereum/Precompiled.cpp b/libethereum/Precompiled.cpp index 0e80949fe..ab0d4d83f 100644 --- a/libethereum/Precompiled.cpp +++ b/libethereum/Precompiled.cpp @@ -31,7 +31,10 @@ using namespace std; using namespace dev; using namespace dev::eth; -static bytes ecrecoverCode(bytesConstRef _in) +namespace +{ + +void ecrecoverCode(bytesConstRef _in, bytesRef _out) { struct inType { @@ -44,47 +47,48 @@ static bytes ecrecoverCode(bytesConstRef _in) memcpy(&in, _in.data(), min(_in.size(), sizeof(in))); h256 ret; - - if ((u256)in.v > 28) - return ret.asBytes(); - SignatureStruct sig(in.r, in.s, (byte)((int)(u256)in.v - 27)); - if (!sig.isValid()) - return ret.asBytes(); - - try + if ((u256)in.v <= 28) { - ret = dev::sha3(recover(sig, in.hash)); + SignatureStruct sig(in.r, in.s, (byte)((int)(u256)in.v - 27)); + if (sig.isValid()) + { + try + { + ret = sha3(recover(sig, in.hash)); + } + catch (...) {} + } } - catch (...) {} memset(ret.data(), 0, 12); - return ret.asBytes(); + ret.ref().copyTo(_out); } -static bytes sha256Code(bytesConstRef _in) +void sha256Code(bytesConstRef _in, bytesRef _out) { - return sha256(_in).asBytes(); + sha256(_in).ref().copyTo(_out); } -static bytes ripemd160Code(bytesConstRef _in) +void ripemd160Code(bytesConstRef _in, bytesRef _out) { - return h256(ripemd160(_in), h256::AlignRight).asBytes(); + h256(ripemd160(_in), h256::AlignRight).ref().copyTo(_out); } -static bytes identityCode(bytesConstRef _in) +void identityCode(bytesConstRef _in, bytesRef _out) { - return _in.toBytes(); + _in.copyTo(_out); } -static const std::unordered_map c_precompiled = -{ - { 1, { [](bytesConstRef) -> bigint { return c_ecrecoverGas; }, ecrecoverCode }}, - { 2, { [](bytesConstRef i) -> bigint { return c_sha256Gas + (i.size() + 31) / 32 * c_sha256WordGas; }, sha256Code }}, - { 3, { [](bytesConstRef i) -> bigint { return c_ripemd160Gas + (i.size() + 31) / 32 * c_ripemd160WordGas; }, ripemd160Code }}, - { 4, { [](bytesConstRef i) -> bigint { return c_identityGas + (i.size() + 31) / 32 * c_identityWordGas; }, identityCode }} -}; +} std::unordered_map const& dev::eth::precompiled() { - return c_precompiled; + static const std::unordered_map s_precompiled = + { + { 1, { [](bytesConstRef) -> bigint { return c_ecrecoverGas; }, ecrecoverCode }}, + { 2, { [](bytesConstRef i) -> bigint { return c_sha256Gas + (i.size() + 31) / 32 * c_sha256WordGas; }, sha256Code }}, + { 3, { [](bytesConstRef i) -> bigint { return c_ripemd160Gas + (i.size() + 31) / 32 * c_ripemd160WordGas; }, ripemd160Code }}, + { 4, { [](bytesConstRef i) -> bigint { return c_identityGas + (i.size() + 31) / 32 * c_identityWordGas; }, identityCode }} + }; + return s_precompiled; } diff --git a/libethereum/Precompiled.h b/libethereum/Precompiled.h index bded27386..b50e51ecd 100644 --- a/libethereum/Precompiled.h +++ b/libethereum/Precompiled.h @@ -34,7 +34,7 @@ namespace eth struct PrecompiledAddress { std::function gas; - std::function exec; + std::function exec; }; /// Info on precompiled contract accounts baked into the protocol. From f042abc3e71af41bd4ce6ca269ceb099ebf9a714 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pawe=C5=82=20Bylica?= Date: Thu, 28 May 2015 10:29:16 +0200 Subject: [PATCH 04/20] Kill steps limit option in VM. --- evmjit/libevmjit-cpp/JitVM.cpp | 4 ++-- evmjit/libevmjit-cpp/JitVM.h | 2 +- libethereum/Executive.cpp | 4 ---- libevm/SmartVM.cpp | 4 ++-- libevm/SmartVM.h | 2 +- libevm/VM.cpp | 9 +++------ libevm/VM.h | 2 +- libevm/VMFace.h | 11 +++++------ 8 files changed, 15 insertions(+), 23 deletions(-) diff --git a/evmjit/libevmjit-cpp/JitVM.cpp b/evmjit/libevmjit-cpp/JitVM.cpp index ac8df545f..0d6a6e00a 100644 --- a/evmjit/libevmjit-cpp/JitVM.cpp +++ b/evmjit/libevmjit-cpp/JitVM.cpp @@ -18,7 +18,7 @@ namespace eth extern "C" void env_sload(); // fake declaration for linker symbol stripping workaround, see a call below -bytesConstRef JitVM::execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp, uint64_t _step) +bytesConstRef JitVM::execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp) { using namespace jit; @@ -33,7 +33,7 @@ bytesConstRef JitVM::execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _on { cwarn << "Execution rejected by EVM JIT (gas limit: " << io_gas << "), executing with interpreter"; m_fallbackVM = VMFactory::create(VMKind::Interpreter); - return m_fallbackVM->execImpl(io_gas, _ext, _onOp, _step); + return m_fallbackVM->execImpl(io_gas, _ext, _onOp); } m_data.gas = static_cast(io_gas); diff --git a/evmjit/libevmjit-cpp/JitVM.h b/evmjit/libevmjit-cpp/JitVM.h index 03c3c8880..e97abd83b 100644 --- a/evmjit/libevmjit-cpp/JitVM.h +++ b/evmjit/libevmjit-cpp/JitVM.h @@ -11,7 +11,7 @@ namespace eth class JitVM: public VMFace { public: - virtual bytesConstRef execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp, uint64_t _steps) override final; + virtual bytesConstRef execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp) override final; private: jit::RuntimeData m_data; diff --git a/libethereum/Executive.cpp b/libethereum/Executive.cpp index e33195631..f3e45cbef 100644 --- a/libethereum/Executive.cpp +++ b/libethereum/Executive.cpp @@ -246,10 +246,6 @@ bool Executive::go(OnOpFunc const& _onOp) vm->exec(m_gas, *m_ext, m_outRef, _onOp); // take only expected output } } - catch (StepsDone const&) - { - return false; - } catch (VMException const& _e) { clog(StateSafeExceptions) << "Safe VM Exception. " << diagnostic_information(_e); diff --git a/libevm/SmartVM.cpp b/libevm/SmartVM.cpp index 27263fc2d..12f2f7078 100644 --- a/libevm/SmartVM.cpp +++ b/libevm/SmartVM.cpp @@ -41,7 +41,7 @@ namespace } } -bytesConstRef SmartVM::execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp, uint64_t _steps) +bytesConstRef SmartVM::execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp) { auto codeHash = sha3(_ext.code); auto vmKind = VMKind::Interpreter; // default VM @@ -68,7 +68,7 @@ bytesConstRef SmartVM::execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _ // TODO: Selected VM must be kept only because it returns reference to its internal memory. // VM implementations should be stateless, without escaping memory reference. m_selectedVM = VMFactory::create(vmKind); - return m_selectedVM->execImpl(io_gas, _ext, _onOp, _steps); + return m_selectedVM->execImpl(io_gas, _ext, _onOp); } } diff --git a/libevm/SmartVM.h b/libevm/SmartVM.h index 73d515f29..fce7be21e 100644 --- a/libevm/SmartVM.h +++ b/libevm/SmartVM.h @@ -31,7 +31,7 @@ namespace eth class SmartVM: public VMFace { public: - virtual bytesConstRef execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp, uint64_t _steps) override final; + virtual bytesConstRef execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp) override final; private: std::unique_ptr m_selectedVM; diff --git a/libevm/VM.cpp b/libevm/VM.cpp index 3a4124d67..dedb30220 100644 --- a/libevm/VM.cpp +++ b/libevm/VM.cpp @@ -45,7 +45,7 @@ static array metrics() return s_ret; } -bytesConstRef VM::execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp, uint64_t _steps) +bytesConstRef VM::execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp) { // Reset leftovers from possible previous run m_curPC = 0; @@ -73,8 +73,7 @@ bytesConstRef VM::execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp, i += _ext.code[i] - (unsigned)Instruction::PUSH1 + 1; } u256 nextPC = m_curPC + 1; - auto osteps = _steps; - for (bool stopped = false; !stopped && _steps--; m_curPC = nextPC, nextPC = m_curPC + 1) + for (uint64_t steps = 0; true; m_curPC = nextPC, nextPC = m_curPC + 1, ++steps) { // INSTRUCTION... Instruction inst = (Instruction)_ext.getCode(m_curPC); @@ -97,7 +96,7 @@ bytesConstRef VM::execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp, auto onOperation = [&]() { if (_onOp) - _onOp(osteps - _steps - 1, inst, newTempSize > m_temp.size() ? (newTempSize - m_temp.size()) / 32 : bigint(0), runGas, io_gas, this, &_ext); + _onOp(steps, inst, newTempSize > m_temp.size() ? (newTempSize - m_temp.size()) / 32 : bigint(0), runGas, io_gas, this, &_ext); }; switch (inst) @@ -670,7 +669,5 @@ bytesConstRef VM::execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp, } } - if (_steps == (uint64_t)-1) - BOOST_THROW_EXCEPTION(StepsDone()); return bytesConstRef(); } diff --git a/libevm/VM.h b/libevm/VM.h index da3a843ed..d87f7e353 100644 --- a/libevm/VM.h +++ b/libevm/VM.h @@ -52,7 +52,7 @@ inline u256 fromAddress(Address _a) class VM: public VMFace { public: - virtual bytesConstRef execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp, uint64_t _steps) override final; + virtual bytesConstRef execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp) override final; u256 curPC() const { return m_curPC; } diff --git a/libevm/VMFace.h b/libevm/VMFace.h index 7a344a072..2a9ed808e 100644 --- a/libevm/VMFace.h +++ b/libevm/VMFace.h @@ -26,7 +26,6 @@ namespace eth { struct VMException: virtual Exception {}; -struct StepsDone: virtual VMException {}; struct BreakPointHit: virtual VMException {}; struct BadInstruction: virtual VMException {}; struct BadJumpDestination: virtual VMException {}; @@ -46,19 +45,19 @@ public: /// Execute EVM code by VM. /// /// @param _out Expected output - void exec(u256& io_gas, ExtVMFace& _ext, bytesRef _out, OnOpFunc const& _onOp = {}, uint64_t _steps = (uint64_t)-1) + void exec(u256& io_gas, ExtVMFace& _ext, bytesRef _out, OnOpFunc const& _onOp = {}) { - execImpl(io_gas, _ext, _onOp, _steps).copyTo(_out); + execImpl(io_gas, _ext, _onOp).copyTo(_out); } /// The same as above but returns a copy of full output. - bytes exec(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp = {}, uint64_t _steps = (uint64_t)-1) + bytes exec(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp = {}) { - return execImpl(io_gas, _ext, _onOp, _steps).toVector(); + return execImpl(io_gas, _ext, _onOp).toVector(); } /// VM implementation - virtual bytesConstRef execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp, uint64_t _steps) = 0; + virtual bytesConstRef execImpl(u256& io_gas, ExtVMFace& _ext, OnOpFunc const& _onOp) = 0; }; } From a10f7d5dd2ac282263f8cb0c5e6b9756525a2f42 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pawe=C5=82=20Bylica?= Date: Thu, 28 May 2015 11:56:28 +0200 Subject: [PATCH 05/20] Rename Executive::setResultRef -> collectResult. --- libethereum/Client.cpp | 2 +- libethereum/Executive.h | 2 +- libethereum/State.cpp | 2 +- mix/MixClient.cpp | 2 +- test/libsolidity/solidityExecutionFramework.h | 2 +- 5 files changed, 5 insertions(+), 5 deletions(-) diff --git a/libethereum/Client.cpp b/libethereum/Client.cpp index d9813c725..14b50e9ba 100644 --- a/libethereum/Client.cpp +++ b/libethereum/Client.cpp @@ -440,7 +440,7 @@ ExecutionResult Client::call(Address _dest, bytes const& _data, u256 _gas, u256 temp = m_postMine; temp.addBalance(_from, _value + _gasPrice * _gas); Executive e(temp, LastHashes(), 0); - e.setResultRef(ret); + e.collectResult(ret); if (!e.call(_dest, _from, _value, _gasPrice, &_data, _gas)) e.go(); e.finalize(); diff --git a/libethereum/Executive.h b/libethereum/Executive.h index 8b2906b9b..668e4b197 100644 --- a/libethereum/Executive.h +++ b/libethereum/Executive.h @@ -115,7 +115,7 @@ public: bool excepted() const { return m_excepted != TransactionException::None; } /// Collect execution results in the result storage provided. - void setResultRef(ExecutionResult& _res) { m_res = &_res; } + void collectResult(ExecutionResult& _res) { m_res = &_res; } private: State& m_s; ///< The state to which this operation/transaction is applied. diff --git a/libethereum/State.cpp b/libethereum/State.cpp index 277e9a649..d96465dd9 100644 --- a/libethereum/State.cpp +++ b/libethereum/State.cpp @@ -1126,7 +1126,7 @@ ExecutionResult State::execute(LastHashes const& _lh, Transaction const& _t, Per // transaction is bad in any way. Executive e(*this, _lh, 0); ExecutionResult res; - e.setResultRef(res); + e.collectResult(res); e.initialize(_t); // Uncommitting is a non-trivial operation - only do it once we've verified as much of the diff --git a/mix/MixClient.cpp b/mix/MixClient.cpp index 2644a8c65..634cf278c 100644 --- a/mix/MixClient.cpp +++ b/mix/MixClient.cpp @@ -132,7 +132,7 @@ void MixClient::executeTransaction(Transaction const& _t, State& _state, bool _c execState.addBalance(t.sender(), t.gas() * t.gasPrice()); //give it enough balance for gas estimation eth::ExecutionResult er; Executive execution(execState, lastHashes, 0); - execution.setResultRef(er); + execution.collectResult(er); execution.initialize(t); execution.execute(); std::vector machineStates; diff --git a/test/libsolidity/solidityExecutionFramework.h b/test/libsolidity/solidityExecutionFramework.h index ea277421a..8b8d851a2 100644 --- a/test/libsolidity/solidityExecutionFramework.h +++ b/test/libsolidity/solidityExecutionFramework.h @@ -136,7 +136,7 @@ protected: m_state.addBalance(m_sender, _value); // just in case eth::Executive executive(m_state, eth::LastHashes(), 0); eth::ExecutionResult res; - executive.setResultRef(res); + executive.collectResult(res); eth::Transaction t = _isCreation ? eth::Transaction(_value, m_gasPrice, m_gas, _data, 0, KeyPair::create().sec()) : eth::Transaction(_value, m_gasPrice, m_gas, m_contractAddress, _data, 0, KeyPair::create().sec()); bytes transactionRLP = t.rlp(); From 96bf8ac00921f8b321e2bedefe7e7a30fd6a08bf Mon Sep 17 00:00:00 2001 From: Kristoffer Josefsson Date: Fri, 29 May 2015 08:44:06 -0700 Subject: [PATCH 06/20] Added chunked upload to older cards by @sontol. --- libethash-cl/ethash_cl_miner.cpp | 386 +++++++++++++++---------- libethash-cl/ethash_cl_miner.h | 6 +- libethash-cl/ethash_cl_miner_kernel.cl | 142 ++++++++- 3 files changed, 378 insertions(+), 156 deletions(-) diff --git a/libethash-cl/ethash_cl_miner.cpp b/libethash-cl/ethash_cl_miner.cpp index be17ba449..be4e7532c 100644 --- a/libethash-cl/ethash_cl_miner.cpp +++ b/libethash-cl/ethash_cl_miner.cpp @@ -24,6 +24,7 @@ #include #include +#include #include #include #include @@ -42,9 +43,13 @@ #define CL_MEM_HOST_READ_ONLY 0 #endif +//#define CHUNKS + #undef min #undef max +//#define CHUNKS + using namespace std; static void add_definition(std::string& source, char const* id, unsigned value) @@ -131,99 +136,147 @@ void ethash_cl_miner::finish() bool ethash_cl_miner::init(uint8_t const* _dag, uint64_t _dagSize, unsigned workgroup_size, unsigned _platformId, unsigned _deviceId) { // get all platforms - std::vector platforms; - cl::Platform::get(&platforms); - if (platforms.empty()) - { - cout << "No OpenCL platforms found." << endl; - return false; - } + try { + std::vector platforms; + cl::Platform::get(&platforms); + if (platforms.empty()) + { + cout << "No OpenCL platforms found." << endl; + return false; + } - // use selected platform + // use selected platform - _platformId = std::min(_platformId, platforms.size() - 1); + _platformId = std::min(_platformId, platforms.size() - 1); - cout << "Using platform: " << platforms[_platformId].getInfo().c_str() << endl; + cout << "Using platform: " << platforms[_platformId].getInfo().c_str() << endl; - // get GPU device of the default platform - std::vector devices; - platforms[_platformId].getDevices(CL_DEVICE_TYPE_ALL, &devices); - if (devices.empty()) - { - cout << "No OpenCL devices found." << endl; - return false; - } + // get GPU device of the default platform + std::vector devices; + platforms[_platformId].getDevices(CL_DEVICE_TYPE_ALL, &devices); + if (devices.empty()) + { + cout << "No OpenCL devices found." << endl; + return false; + } - // use selected device - cl::Device& device = devices[std::min(_deviceId, devices.size() - 1)]; - std::string device_version = device.getInfo(); - cout << "Using device: " << device.getInfo().c_str() << "(" << device_version.c_str() << ")" << endl; + // use selected device + cl::Device& device = devices[std::min(_deviceId, devices.size() - 1)]; + std::string device_version = device.getInfo(); + cout << "Using device: " << device.getInfo().c_str() << "(" << device_version.c_str() << ")" << endl; - if (strncmp("OpenCL 1.0", device_version.c_str(), 10) == 0) - { - cout << "OpenCL 1.0 is not supported." << endl; - return false; - } - if (strncmp("OpenCL 1.1", device_version.c_str(), 10) == 0) - m_opencl_1_1 = true; - - // create context - m_context = cl::Context(std::vector(&device, &device + 1)); - m_queue = cl::CommandQueue(m_context, device); - - // use requested workgroup size, but we require multiple of 8 - m_workgroup_size = ((workgroup_size + 7) / 8) * 8; - - // patch source code - std::string code(ETHASH_CL_MINER_KERNEL, ETHASH_CL_MINER_KERNEL + ETHASH_CL_MINER_KERNEL_SIZE); - add_definition(code, "GROUP_SIZE", m_workgroup_size); - add_definition(code, "DAG_SIZE", (unsigned)(_dagSize / ETHASH_MIX_BYTES)); - add_definition(code, "ACCESSES", ETHASH_ACCESSES); - add_definition(code, "MAX_OUTPUTS", c_max_search_results); - //debugf("%s", code.c_str()); - - // create miner OpenCL program - cl::Program::Sources sources; - sources.push_back({code.c_str(), code.size()}); - - cl::Program program(m_context, sources); - try - { - program.build({device}); + if (strncmp("OpenCL 1.0", device_version.c_str(), 10) == 0) + { + cout << "OpenCL 1.0 is not supported." << endl; + return false; + } + if (strncmp("OpenCL 1.1", device_version.c_str(), 10) == 0) + m_opencl_1_1 = true; + + // create context + m_context = cl::Context(std::vector(&device, &device + 1)); + m_queue = cl::CommandQueue(m_context, device); + + // use requested workgroup size, but we require multiple of 8 + m_workgroup_size = ((workgroup_size + 7) / 8) * 8; + + // patch source code + std::ifstream t("ethash_cl_miner_kernel.cl"); + std::string code((std::istreambuf_iterator(t)), + std::istreambuf_iterator()); + add_definition(code, "GROUP_SIZE", m_workgroup_size); + add_definition(code, "DAG_SIZE", (unsigned)(_dagSize / ETHASH_MIX_BYTES)); + add_definition(code, "ACCESSES", ETHASH_ACCESSES); + add_definition(code, "MAX_OUTPUTS", c_max_search_results); + //debugf("%s", code.c_str()); + + // create miner OpenCL program + cl::Program::Sources sources; + sources.push_back({ code.c_str(), code.size() }); + + cl::Program program(m_context, sources); + try + { + program.build({ device }); + cout << "Printing program log" << endl; + cout << program.getBuildInfo(device).c_str(); + } + catch (cl::Error err) + { + cout << program.getBuildInfo(device).c_str(); + return false; + } + #ifdef CHUNKS + cout << "loading ethash_hash_chunks" << endl; + m_hash_kernel = cl::Kernel(program, "ethash_hash_chunks"); + cout << "loading ethash_search_chunks" << endl; + m_search_kernel = cl::Kernel(program, "ethash_search_chunks"); + + #else + cout << "loading ethash_hash" << endl; + m_hash_kernel = cl::Kernel(program, "ethash_hash"); + cout << "loading ethash_search" << endl; + m_search_kernel = cl::Kernel(program, "ethash_search"); + #endif + + // create buffer for dag + #ifdef CHUNKS + for (unsigned i = 0; i < 4; i++){ + + cout << "Creating chunky buffer: " << i << endl; + m_dags[i] = cl::Buffer(m_context, CL_MEM_READ_ONLY, (i == 3) ? (_dagSize - 3 * ((_dagSize >> 9) << 7)) : (_dagSize >> 9) << 7); + } + #else + cout << "Creating one big buffer." << endl; + m_dag = cl::Buffer(m_context, CL_MEM_READ_ONLY, _dagSize); + #endif + + // create buffer for header + cout << "Creating buffer for header." << endl; + m_header = cl::Buffer(m_context, CL_MEM_READ_ONLY, 32); + + #ifdef CHUNKS + void* dag_ptr[4]; + for (unsigned i = 0; i < 4; i++) + { + cout << "Mapping chunk " << i << endl; + dag_ptr[i] = m_queue.enqueueMapBuffer(m_dags[i], true, m_opencl_1_1 ? CL_MAP_WRITE : CL_MAP_WRITE_INVALIDATE_REGION, 0, (i == 3) ? (_dagSize - 3 * ((_dagSize >> 9) << 7)) : (_dagSize >> 9) << 7); + } + for (unsigned i = 0; i < 4; i++) + { + memcpy(dag_ptr[i], (char *)_dag + i*((_dagSize >> 9) << 7), (i == 3) ? (_dagSize - 3 * ((_dagSize >> 9) << 7)) : (_dagSize >> 9) << 7); + m_queue.enqueueUnmapMemObject(m_dags[i], dag_ptr[i]); + } + #else + cout << "Mapping chunk." << endl; + m_queue.enqueueWriteBuffer(m_dag, CL_TRUE, 0, _dagSize, _dag); + #endif + // compute dag on CPU + /*{ + m_queue.enqueueWriteBuffer(m_dag, CL_TRUE, 0, _dagSize, _dag); + + // if this throws then it's because we probably need to subdivide the dag uploads for compatibility + // void* dag_ptr = m_queue.enqueueMapBuffer(m_dag, true, m_opencl_1_1 ? CL_MAP_WRITE : CL_MAP_WRITE_INVALIDATE_REGION, 0, _dagSize); + // memcpying 1GB: horrible... really. horrible. but necessary since we can't mmap *and* gpumap. + // _fillDAG(dag_ptr); + // m_queue.enqueueUnmapMemObject(m_dag, dag_ptr); + }*/ + + // create mining buffers + for (unsigned i = 0; i != c_num_buffers; ++i) + { + cout << "Creating minig buffer " << i <(device).c_str(); - return false; + std::cout << err.what() << "(" << err.err() << ")" << std::endl; } - m_hash_kernel = cl::Kernel(program, "ethash_hash"); - m_search_kernel = cl::Kernel(program, "ethash_search"); - - // create buffer for dag - m_dag = cl::Buffer(m_context, CL_MEM_READ_ONLY, _dagSize); - - // create buffer for header - m_header = cl::Buffer(m_context, CL_MEM_READ_ONLY, 32); - // compute dag on CPU - try { - m_queue.enqueueWriteBuffer(m_dag, CL_TRUE, 0, _dagSize, _dag); - } - catch (...) - { - // didn't work. shitty driver. try allocating in CPU RAM and manually memcpying it. - void* dag_ptr = m_queue.enqueueMapBuffer(m_dag, true, m_opencl_1_1 ? CL_MAP_WRITE : CL_MAP_WRITE_INVALIDATE_REGION, 0, _dagSize); - memcpy(dag_ptr, _dag, _dagSize); - m_queue.enqueueUnmapMemObject(m_dag, dag_ptr); - } - - // create mining buffers - for (unsigned i = 0; i != c_num_buffers; ++i) - { - m_hash_buf[i] = cl::Buffer(m_context, CL_MEM_WRITE_ONLY | (!m_opencl_1_1 ? CL_MEM_HOST_READ_ONLY : 0), 32*c_hash_batch_size); - m_search_buf[i] = cl::Buffer(m_context, CL_MEM_WRITE_ONLY, (c_max_search_results + 1) * sizeof(uint32_t)); - } - return true; + return true; } void ethash_cl_miner::hash(uint8_t* ret, uint8_t const* header, uint64_t nonce, unsigned count) @@ -248,10 +301,22 @@ void ethash_cl_miner::hash(uint8_t* ret, uint8_t const* header, uint64_t nonce, uint isolate ) */ + #ifdef CHUNKS + cout << "Setting chunk hash arguments." << endl; + m_hash_kernel.setArg(1, m_header); + m_hash_kernel.setArg(2, m_dags[0]); + m_hash_kernel.setArg(3, m_dags[1]); + m_hash_kernel.setArg(4, m_dags[2]); + m_hash_kernel.setArg(5, m_dags[3]); + m_hash_kernel.setArg(6, nonce); + m_hash_kernel.setArg(7, ~0u); // have to pass this to stop the compile unrolling the loop + #else + cout << "Setting hash arguments." << endl; m_hash_kernel.setArg(1, m_header); m_hash_kernel.setArg(2, m_dag); m_hash_kernel.setArg(3, nonce); m_hash_kernel.setArg(4, ~0u); // have to pass this to stop the compile unrolling the loop + #endif unsigned buf = 0; for (unsigned i = 0; i < count || !pending.empty(); ) @@ -297,95 +362,120 @@ void ethash_cl_miner::hash(uint8_t* ret, uint8_t const* header, uint64_t nonce, void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook& hook) { - struct pending_batch - { - uint64_t start_nonce; - unsigned buf; - }; - std::queue pending; + try { + struct pending_batch + { + uint64_t start_nonce; + unsigned buf; + }; + std::queue pending; - static uint32_t const c_zero = 0; + static uint32_t const c_zero = 0; - // update header constant buffer - m_queue.enqueueWriteBuffer(m_header, false, 0, 32, header); - for (unsigned i = 0; i != c_num_buffers; ++i) - m_queue.enqueueWriteBuffer(m_search_buf[i], false, 0, 4, &c_zero); + // update header constant buffer + m_queue.enqueueWriteBuffer(m_header, false, 0, 32, header); + for (unsigned i = 0; i != c_num_buffers; ++i) + m_queue.enqueueWriteBuffer(m_search_buf[i], false, 0, 4, &c_zero); #if CL_VERSION_1_2 && 0 - cl::Event pre_return_event; - if (!m_opencl_1_1) - m_queue.enqueueBarrierWithWaitList(NULL, &pre_return_event); - else + cl::Event pre_return_event; + if (!m_opencl_1_1) + m_queue.enqueueBarrierWithWaitList(NULL, &pre_return_event); + else #endif - m_queue.finish(); + m_queue.finish(); - /* - __kernel void ethash_combined_search( + /* + __kernel void ethash_combined_search( __global hash32_t* g_hashes, // 0 __constant hash32_t const* g_header, // 1 __global hash128_t const* g_dag, // 2 ulong start_nonce, // 3 ulong target, // 4 uint isolate // 5 - ) - */ - m_search_kernel.setArg(1, m_header); - m_search_kernel.setArg(2, m_dag); - - // pass these to stop the compiler unrolling the loops - m_search_kernel.setArg(4, target); - m_search_kernel.setArg(5, ~0u); + ) + */ + #ifdef CHUNKS + cout << "Setting chunk search arguments." << endl; + m_search_kernel.setArg(1, m_header); + m_search_kernel.setArg(2, m_dags[0]); + m_search_kernel.setArg(3, m_dags[1]); + m_search_kernel.setArg(4, m_dags[2]); + m_search_kernel.setArg(5, m_dags[3]); + + // pass these to stop the compiler unrolling the loops + m_search_kernel.setArg(7, target); + m_search_kernel.setArg(8, ~0u); + + #else + cout << "Setting search arguments." << endl; + m_search_kernel.setArg(1, m_header); + m_search_kernel.setArg(2, m_dag); + + // pass these to stop the compiler unrolling the loops + m_search_kernel.setArg(4, target); + m_search_kernel.setArg(5, ~0u); + #endif + + + + unsigned buf = 0; + std::random_device engine; + uint64_t start_nonce = std::uniform_int_distribution()(engine); + for (;; start_nonce += c_search_batch_size) + { + // supply output buffer to kernel + m_search_kernel.setArg(0, m_search_buf[buf]); + #ifdef CHUNKS + m_search_kernel.setArg(6, start_nonce); + #else + m_search_kernel.setArg(3, start_nonce); + #endif + // execute it! + m_queue.enqueueNDRangeKernel(m_search_kernel, cl::NullRange, c_search_batch_size, m_workgroup_size); - unsigned buf = 0; - std::random_device engine; - uint64_t start_nonce = std::uniform_int_distribution()(engine); - for (; ; start_nonce += c_search_batch_size) - { - // supply output buffer to kernel - m_search_kernel.setArg(0, m_search_buf[buf]); - m_search_kernel.setArg(3, start_nonce); + pending.push({ start_nonce, buf }); + buf = (buf + 1) % c_num_buffers; - // execute it! - m_queue.enqueueNDRangeKernel(m_search_kernel, cl::NullRange, c_search_batch_size, m_workgroup_size); - - pending.push({start_nonce, buf}); - buf = (buf + 1) % c_num_buffers; + // read results + if (pending.size() == c_num_buffers) + { + pending_batch const& batch = pending.front(); - // read results - if (pending.size() == c_num_buffers) - { - pending_batch const& batch = pending.front(); + // could use pinned host pointer instead + uint32_t* results = (uint32_t*)m_queue.enqueueMapBuffer(m_search_buf[batch.buf], true, CL_MAP_READ, 0, (1 + c_max_search_results) * sizeof(uint32_t)); + unsigned num_found = std::min(results[0], c_max_search_results); - // could use pinned host pointer instead - uint32_t* results = (uint32_t*)m_queue.enqueueMapBuffer(m_search_buf[batch.buf], true, CL_MAP_READ, 0, (1+c_max_search_results) * sizeof(uint32_t)); - unsigned num_found = std::min(results[0], c_max_search_results); + uint64_t nonces[c_max_search_results]; + for (unsigned i = 0; i != num_found; ++i) + { + nonces[i] = batch.start_nonce + results[i + 1]; + } - uint64_t nonces[c_max_search_results]; - for (unsigned i = 0; i != num_found; ++i) - { - nonces[i] = batch.start_nonce + results[i+1]; - } + m_queue.enqueueUnmapMemObject(m_search_buf[batch.buf], results); - m_queue.enqueueUnmapMemObject(m_search_buf[batch.buf], results); - - bool exit = num_found && hook.found(nonces, num_found); - exit |= hook.searched(batch.start_nonce, c_search_batch_size); // always report searched before exit - if (exit) - break; + bool exit = num_found && hook.found(nonces, num_found); + exit |= hook.searched(batch.start_nonce, c_search_batch_size); // always report searched before exit + if (exit) + break; - // reset search buffer if we're still going - if (num_found) - m_queue.enqueueWriteBuffer(m_search_buf[batch.buf], true, 0, 4, &c_zero); + // reset search buffer if we're still going + if (num_found) + m_queue.enqueueWriteBuffer(m_search_buf[batch.buf], true, 0, 4, &c_zero); - pending.pop(); + pending.pop(); + } } - } - // not safe to return until this is ready + // not safe to return until this is ready #if CL_VERSION_1_2 && 0 - if (!m_opencl_1_1) - pre_return_event.wait(); + if (!m_opencl_1_1) + pre_return_event.wait(); #endif + } + catch (cl::Error err) + { + std::cout << err.what() << "(" << err.err() << ")" << std::endl; + } } - diff --git a/libethash-cl/ethash_cl_miner.h b/libethash-cl/ethash_cl_miner.h index 43bfa2336..9c97f2aa4 100644 --- a/libethash-cl/ethash_cl_miner.h +++ b/libethash-cl/ethash_cl_miner.h @@ -41,6 +41,9 @@ public: void hash(uint8_t* ret, uint8_t const* header, uint64_t nonce, unsigned count); void search(uint8_t const* header, uint64_t target, search_hook& hook); + void hash_chunk(uint8_t* ret, uint8_t const* header, uint64_t nonce, unsigned count); + void search_chunk(uint8_t const* header, uint64_t target, search_hook& hook); + private: enum { c_max_search_results = 63, c_num_buffers = 2, c_hash_batch_size = 1024, c_search_batch_size = 1024*256 }; @@ -49,9 +52,10 @@ private: cl::Kernel m_hash_kernel; cl::Kernel m_search_kernel; cl::Buffer m_dag; + cl::Buffer m_dags[4]; cl::Buffer m_header; cl::Buffer m_hash_buf[c_num_buffers]; cl::Buffer m_search_buf[c_num_buffers]; unsigned m_workgroup_size; bool m_opencl_1_1; -}; +}; \ No newline at end of file diff --git a/libethash-cl/ethash_cl_miner_kernel.cl b/libethash-cl/ethash_cl_miner_kernel.cl index 3c8b9dc92..8567bb164 100644 --- a/libethash-cl/ethash_cl_miner_kernel.cl +++ b/libethash-cl/ethash_cl_miner_kernel.cl @@ -179,13 +179,13 @@ void keccak_f1600_no_absorb(ulong* a, uint in_size, uint out_size, uint isolate) // much we try and help the compiler save VGPRs because it seems to throw // that information away, hence the implementation of keccak here // doesn't bother. - if (isolate) + if (isolate) { keccak_f1600_round((uint2*)a, r++, 25); } } while (r < 23); - + // final round optimised for digest size keccak_f1600_round((uint2*)a, r++, out_size); } @@ -232,7 +232,7 @@ hash64_t init_hash(__constant hash32_t const* header, ulong nonce, uint isolate) hash64_t init; uint const init_size = countof(init.ulongs); uint const hash_size = countof(header->ulongs); - + // sha3_512(header .. nonce) ulong state[25]; copy(state, header->ulongs, hash_size); @@ -243,6 +243,40 @@ hash64_t init_hash(__constant hash32_t const* header, ulong nonce, uint isolate) return init; } +uint inner_loop_chunks(uint4 init, uint thread_id, __local uint* share, __global hash128_t const* g_dag, __global hash128_t const* g_dag1, __global hash128_t const* g_dag2, __global hash128_t const* g_dag3, uint isolate) +{ + uint4 mix = init; + + // share init0 + if (thread_id == 0) + *share = mix.x; + barrier(CLK_LOCAL_MEM_FENCE); + uint init0 = *share; + + uint a = 0; + do + { + bool update_share = thread_id == (a/4) % THREADS_PER_HASH; + + #pragma unroll + for (uint i = 0; i != 4; ++i) + { + if (update_share) + { + uint m[4] = { mix.x, mix.y, mix.z, mix.w }; + *share = fnv(init0 ^ (a+i), m[i]) % DAG_SIZE; + } + barrier(CLK_LOCAL_MEM_FENCE); + + mix = fnv4(mix, *share>=3 * DAG_SIZE / 4 ? g_dag3[*share - 3 * DAG_SIZE / 4].uint4s[thread_id] : *share>=DAG_SIZE / 2 ? g_dag2[*share - DAG_SIZE / 2].uint4s[thread_id] : *share>=DAG_SIZE / 4 ? g_dag1[*share - DAG_SIZE / 4].uint4s[thread_id]:g_dag[*share].uint4s[thread_id]); + } + } while ((a += 4) != (ACCESSES & isolate)); + + return fnv_reduce(mix); +} + + + uint inner_loop(uint4 init, uint thread_id, __local uint* share, __global hash128_t const* g_dag, uint isolate) { uint4 mix = init; @@ -276,6 +310,7 @@ uint inner_loop(uint4 init, uint thread_id, __local uint* share, __global hash12 return fnv_reduce(mix); } + hash32_t final_hash(hash64_t const* init, hash32_t const* mix, uint isolate) { ulong state[25]; @@ -309,7 +344,7 @@ hash32_t compute_hash_simple( { mix.uint4s[i] = init.uint4s[i % countof(init.uint4s)]; } - + uint mix_val = mix.uints[0]; uint init0 = mix.uints[0]; uint a = 0; @@ -333,7 +368,7 @@ hash32_t compute_hash_simple( { fnv_mix.uints[i] = fnv_reduce(mix.uint4s[i]); } - + return final_hash(&init, &fnv_mix, isolate); } @@ -347,6 +382,7 @@ typedef union hash32_t mix; } compute_hash_share; + hash32_t compute_hash( __local compute_hash_share* share, __constant hash32_t const* g_header, @@ -390,6 +426,53 @@ hash32_t compute_hash( return final_hash(&init, &mix, isolate); } + +hash32_t compute_hash_chunks( + __local compute_hash_share* share, + __constant hash32_t const* g_header, + __global hash128_t const* g_dag, + __global hash128_t const* g_dag1, + __global hash128_t const* g_dag2, + __global hash128_t const* g_dag3, + ulong nonce, + uint isolate + ) +{ + uint const gid = get_global_id(0); + + // Compute one init hash per work item. + hash64_t init = init_hash(g_header, nonce, isolate); + + // Threads work together in this phase in groups of 8. + uint const thread_id = gid % THREADS_PER_HASH; + uint const hash_id = (gid % GROUP_SIZE) / THREADS_PER_HASH; + + hash32_t mix; + uint i = 0; + do + { + // share init with other threads + if (i == thread_id) + share[hash_id].init = init; + barrier(CLK_LOCAL_MEM_FENCE); + + uint4 thread_init = share[hash_id].init.uint4s[thread_id % (64 / sizeof(uint4))]; + barrier(CLK_LOCAL_MEM_FENCE); + + uint thread_mix = inner_loop_chunks(thread_init, thread_id, share[hash_id].mix.uints, g_dag, g_dag1, g_dag2, g_dag3, isolate); + + share[hash_id].mix.uints[thread_id] = thread_mix; + barrier(CLK_LOCAL_MEM_FENCE); + + if (i == thread_id) + mix = share[hash_id].mix; + barrier(CLK_LOCAL_MEM_FENCE); + } + while (++i != (THREADS_PER_HASH & isolate)); + + return final_hash(&init, &mix, isolate); +} + __attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1))) __kernel void ethash_hash_simple( __global hash32_t* g_hashes, @@ -415,13 +498,15 @@ __kernel void ethash_search_simple( { uint const gid = get_global_id(0); hash32_t hash = compute_hash_simple(g_header, g_dag, start_nonce + gid, isolate); - if (as_ulong(as_uchar8(hash.ulongs[0]).s76543210) < target) + + if (hash.ulongs[countof(hash.ulongs)-1] < target) { - uint slot = min(MAX_OUTPUTS, atomic_inc(&g_output[0]) + 1); + uint slot = min(convert_uint(MAX_OUTPUTS), convert_uint(atomic_inc(&g_output[0]) + 1)); g_output[slot] = gid; } } + __attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1))) __kernel void ethash_hash( __global hash32_t* g_hashes, @@ -458,3 +543,46 @@ __kernel void ethash_search( g_output[slot] = gid; } } + +__attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1))) +__kernel void ethash_hash_chunks( + __global hash32_t* g_hashes, + __constant hash32_t const* g_header, + __global hash128_t const* g_dag, + __global hash128_t const* g_dag1, + __global hash128_t const* g_dag2, + __global hash128_t const* g_dag3, + ulong start_nonce, + uint isolate + ) +{ + __local compute_hash_share share[HASHES_PER_LOOP]; + + uint const gid = get_global_id(0); + g_hashes[gid] = compute_hash_chunks(share, g_header, g_dag, g_dag1, g_dag2, g_dag3,start_nonce + gid, isolate); +} + +__attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1))) +__kernel void ethash_search_chunks( + __global volatile uint* restrict g_output, + __constant hash32_t const* g_header, + __global hash128_t const* g_dag, + __global hash128_t const* g_dag1, + __global hash128_t const* g_dag2, + __global hash128_t const* g_dag3, + ulong start_nonce, + ulong target, + uint isolate + ) +{ + __local compute_hash_share share[HASHES_PER_LOOP]; + + uint const gid = get_global_id(0); + hash32_t hash = compute_hash_chunks(share, g_header, g_dag, g_dag1, g_dag2, g_dag3, start_nonce + gid, isolate); + + if (as_ulong(as_uchar8(hash.ulongs[0]).s76543210) < target) + { + uint slot = min(convert_uint(MAX_OUTPUTS), convert_uint(atomic_inc(&g_output[0]) + 1)); + g_output[slot] = gid; + } +} \ No newline at end of file From aeb49b809135e5db9a28ad20443f23abf25114d9 Mon Sep 17 00:00:00 2001 From: Lefteris Karapetsas Date: Mon, 1 Jun 2015 13:03:02 +0200 Subject: [PATCH 07/20] Style changes --- libethash-cl/ethash_cl_miner.cpp | 75 ++++++++++++-------------------- 1 file changed, 28 insertions(+), 47 deletions(-) diff --git a/libethash-cl/ethash_cl_miner.cpp b/libethash-cl/ethash_cl_miner.cpp index be4e7532c..520a13180 100644 --- a/libethash-cl/ethash_cl_miner.cpp +++ b/libethash-cl/ethash_cl_miner.cpp @@ -43,13 +43,12 @@ #define CL_MEM_HOST_READ_ONLY 0 #endif -//#define CHUNKS +// maybe move to CMakeLists.txt ? +// #define ETHASH_CL_CHUNK_UPLOAD #undef min #undef max -//#define CHUNKS - using namespace std; static void add_definition(std::string& source, char const* id, unsigned value) @@ -136,7 +135,8 @@ void ethash_cl_miner::finish() bool ethash_cl_miner::init(uint8_t const* _dag, uint64_t _dagSize, unsigned workgroup_size, unsigned _platformId, unsigned _deviceId) { // get all platforms - try { + try + { std::vector platforms; cl::Platform::get(&platforms); if (platforms.empty()) @@ -146,7 +146,6 @@ bool ethash_cl_miner::init(uint8_t const* _dag, uint64_t _dagSize, unsigned work } // use selected platform - _platformId = std::min(_platformId, platforms.size() - 1); cout << "Using platform: " << platforms[_platformId].getInfo().c_str() << endl; @@ -206,12 +205,11 @@ bool ethash_cl_miner::init(uint8_t const* _dag, uint64_t _dagSize, unsigned work cout << program.getBuildInfo(device).c_str(); return false; } - #ifdef CHUNKS + #ifdef ETHASH_CL_CHUNK_UPLOAD cout << "loading ethash_hash_chunks" << endl; m_hash_kernel = cl::Kernel(program, "ethash_hash_chunks"); cout << "loading ethash_search_chunks" << endl; m_search_kernel = cl::Kernel(program, "ethash_search_chunks"); - #else cout << "loading ethash_hash" << endl; m_hash_kernel = cl::Kernel(program, "ethash_hash"); @@ -220,10 +218,10 @@ bool ethash_cl_miner::init(uint8_t const* _dag, uint64_t _dagSize, unsigned work #endif // create buffer for dag - #ifdef CHUNKS - for (unsigned i = 0; i < 4; i++){ - - cout << "Creating chunky buffer: " << i << endl; + #ifdef ETHASH_CL_CHUNK_UPLOAD + for (unsigned i = 0; i < 4; i++) + { + cout << "Creating chunky buffer: " << i << endl; m_dags[i] = cl::Buffer(m_context, CL_MEM_READ_ONLY, (i == 3) ? (_dagSize - 3 * ((_dagSize >> 9) << 7)) : (_dagSize >> 9) << 7); } #else @@ -235,7 +233,7 @@ bool ethash_cl_miner::init(uint8_t const* _dag, uint64_t _dagSize, unsigned work cout << "Creating buffer for header." << endl; m_header = cl::Buffer(m_context, CL_MEM_READ_ONLY, 32); - #ifdef CHUNKS + #ifdef ETHASH_CL_CHUNK_UPLOAD void* dag_ptr[4]; for (unsigned i = 0; i < 4; i++) { @@ -247,36 +245,25 @@ bool ethash_cl_miner::init(uint8_t const* _dag, uint64_t _dagSize, unsigned work memcpy(dag_ptr[i], (char *)_dag + i*((_dagSize >> 9) << 7), (i == 3) ? (_dagSize - 3 * ((_dagSize >> 9) << 7)) : (_dagSize >> 9) << 7); m_queue.enqueueUnmapMemObject(m_dags[i], dag_ptr[i]); } - #else - cout << "Mapping chunk." << endl; - m_queue.enqueueWriteBuffer(m_dag, CL_TRUE, 0, _dagSize, _dag); - #endif - // compute dag on CPU - /*{ - m_queue.enqueueWriteBuffer(m_dag, CL_TRUE, 0, _dagSize, _dag); - - // if this throws then it's because we probably need to subdivide the dag uploads for compatibility - // void* dag_ptr = m_queue.enqueueMapBuffer(m_dag, true, m_opencl_1_1 ? CL_MAP_WRITE : CL_MAP_WRITE_INVALIDATE_REGION, 0, _dagSize); - // memcpying 1GB: horrible... really. horrible. but necessary since we can't mmap *and* gpumap. - // _fillDAG(dag_ptr); - // m_queue.enqueueUnmapMemObject(m_dag, dag_ptr); - }*/ + #else + cout << "Mapping chunk." << endl; + m_queue.enqueueWriteBuffer(m_dag, CL_TRUE, 0, _dagSize, _dag); + #endif // create mining buffers for (unsigned i = 0; i != c_num_buffers; ++i) { - cout << "Creating minig buffer " << i < pending; - + // update header constant buffer m_queue.enqueueWriteBuffer(m_header, true, 0, 32, header); @@ -301,8 +288,8 @@ void ethash_cl_miner::hash(uint8_t* ret, uint8_t const* header, uint64_t nonce, uint isolate ) */ - #ifdef CHUNKS - cout << "Setting chunk hash arguments." << endl; + #ifdef ETHASH_CL_CHUNK_UPLOAD + cout << "Setting chunk hash arguments." << endl; m_hash_kernel.setArg(1, m_header); m_hash_kernel.setArg(2, m_dags[0]); m_hash_kernel.setArg(3, m_dags[1]); @@ -316,7 +303,7 @@ void ethash_cl_miner::hash(uint8_t* ret, uint8_t const* header, uint64_t nonce, m_hash_kernel.setArg(2, m_dag); m_hash_kernel.setArg(3, nonce); m_hash_kernel.setArg(4, ~0u); // have to pass this to stop the compile unrolling the loop - #endif + #endif unsigned buf = 0; for (unsigned i = 0; i < count || !pending.empty(); ) @@ -336,9 +323,9 @@ void ethash_cl_miner::hash(uint8_t* ret, uint8_t const* header, uint64_t nonce, cl::NullRange, cl::NDRange(batch_count), cl::NDRange(m_workgroup_size) - ); + ); m_queue.flush(); - + pending.push({i, this_count, buf}); i += this_count; buf = (buf + 1) % c_num_buffers; @@ -348,12 +335,10 @@ void ethash_cl_miner::hash(uint8_t* ret, uint8_t const* header, uint64_t nonce, if (i == count || pending.size() == c_num_buffers) { pending_batch const& batch = pending.front(); - // could use pinned host pointer instead, but this path isn't that important. uint8_t* hashes = (uint8_t*)m_queue.enqueueMapBuffer(m_hash_buf[batch.buf], true, CL_MAP_READ, 0, batch.count * ETHASH_BYTES); memcpy(ret + batch.base*ETHASH_BYTES, hashes, batch.count*ETHASH_BYTES); m_queue.enqueueUnmapMemObject(m_hash_buf[batch.buf], hashes); - pending.pop(); } } @@ -362,7 +347,8 @@ void ethash_cl_miner::hash(uint8_t* ret, uint8_t const* header, uint64_t nonce, void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook& hook) { - try { + try + { struct pending_batch { uint64_t start_nonce; @@ -395,7 +381,7 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook uint isolate // 5 ) */ - #ifdef CHUNKS + #ifdef ETHASH_CL_CHUNK_UPLOAD cout << "Setting chunk search arguments." << endl; m_search_kernel.setArg(1, m_header); m_search_kernel.setArg(2, m_dags[0]); @@ -407,8 +393,8 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook m_search_kernel.setArg(7, target); m_search_kernel.setArg(8, ~0u); - #else - cout << "Setting search arguments." << endl; + #else + cout << "Setting search arguments." << endl; m_search_kernel.setArg(1, m_header); m_search_kernel.setArg(2, m_dag); @@ -417,8 +403,6 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook m_search_kernel.setArg(5, ~0u); #endif - - unsigned buf = 0; std::random_device engine; uint64_t start_nonce = std::uniform_int_distribution()(engine); @@ -426,7 +410,7 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook { // supply output buffer to kernel m_search_kernel.setArg(0, m_search_buf[buf]); - #ifdef CHUNKS + #ifdef ETHASH_CL_CHUNK_UPLOAD m_search_kernel.setArg(6, start_nonce); #else m_search_kernel.setArg(3, start_nonce); @@ -449,12 +433,9 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook uint64_t nonces[c_max_search_results]; for (unsigned i = 0; i != num_found; ++i) - { nonces[i] = batch.start_nonce + results[i + 1]; - } m_queue.enqueueUnmapMemObject(m_search_buf[batch.buf], results); - bool exit = num_found && hook.found(nonces, num_found); exit |= hook.searched(batch.start_nonce, c_search_batch_size); // always report searched before exit if (exit) From 587209cf5dd86702de000696940b7075d393bfb0 Mon Sep 17 00:00:00 2001 From: Lefteris Karapetsas Date: Mon, 1 Jun 2015 14:49:58 +0200 Subject: [PATCH 08/20] GPU DAG Chunks is now dynamic argument By providing the --use-chunks argument dagChunks is set to 4. Default is 1 big chunk. Future improvement could be to provide arbitrary number of chunks. --- ethminer/MinerAux.h | 6 ++ libethash-cl/ethash_cl_miner.cpp | 150 +++++++++++++++---------------- libethash-cl/ethash_cl_miner.h | 15 +++- libethcore/Ethash.cpp | 3 +- libethcore/Ethash.h | 2 + 5 files changed, 96 insertions(+), 80 deletions(-) diff --git a/ethminer/MinerAux.h b/ethminer/MinerAux.h index 47fd2e2ae..b6d87e181 100644 --- a/ethminer/MinerAux.h +++ b/ethminer/MinerAux.h @@ -127,6 +127,10 @@ public: cerr << "Bad " << arg << " option: " << argv[i] << endl; throw BadArgument(); } + else if (arg == "--use-chunks") + { + dagChunks = 4; + } else if (arg == "--phone-home" && i + 1 < argc) { string m = argv[++i]; @@ -293,6 +297,7 @@ public: << " --opencl-platform When mining using -G/--opencl use OpenCL platform n (default: 0)." << endl << " --opencl-device When mining using -G/--opencl use OpenCL device n (default: 0)." << endl << " -t, --mining-threads Limit number of CPU/GPU miners to n (default: use everything available on selected platform)" << endl + << " --use-chunks When using GPU mining upload the DAG to the GPU in 4 chunks. " << endl ; } @@ -480,6 +485,7 @@ private: unsigned openclPlatform = 0; unsigned openclDevice = 0; unsigned miningThreads = UINT_MAX; + unsigned dagChunks = 1; /// DAG initialisation param. unsigned initDAG = 0; diff --git a/libethash-cl/ethash_cl_miner.cpp b/libethash-cl/ethash_cl_miner.cpp index 520a13180..5e69df1ee 100644 --- a/libethash-cl/ethash_cl_miner.cpp +++ b/libethash-cl/ethash_cl_miner.cpp @@ -43,9 +43,6 @@ #define CL_MEM_HOST_READ_ONLY 0 #endif -// maybe move to CMakeLists.txt ? -// #define ETHASH_CL_CHUNK_UPLOAD - #undef min #undef max @@ -61,7 +58,7 @@ static void add_definition(std::string& source, char const* id, unsigned value) ethash_cl_miner::search_hook::~search_hook() {} ethash_cl_miner::ethash_cl_miner() -: m_opencl_1_1() +: m_dagChunks(nullptr), m_opencl_1_1() { } @@ -130,10 +127,26 @@ void ethash_cl_miner::finish() { if (m_queue()) m_queue.finish(); + + if (m_dagChunks) + delete [] m_dagChunks; } -bool ethash_cl_miner::init(uint8_t const* _dag, uint64_t _dagSize, unsigned workgroup_size, unsigned _platformId, unsigned _deviceId) +bool ethash_cl_miner::init( + uint8_t const* _dag, + uint64_t _dagSize, + unsigned workgroup_size, + unsigned _platformId, + unsigned _deviceId, + unsigned _dagChunksNum +) { + // for now due to the .cl kernels we can only have either 1 big chunk or 4 chunks + assert(_dagChunksNum == 1 || _dagChunksNum == 4); + // now create the number of chunk buffers + m_dagChunks = new cl::Buffer[_dagChunksNum]; + m_dagChunksNum = _dagChunksNum; + // get all platforms try { @@ -205,50 +218,61 @@ bool ethash_cl_miner::init(uint8_t const* _dag, uint64_t _dagSize, unsigned work cout << program.getBuildInfo(device).c_str(); return false; } - #ifdef ETHASH_CL_CHUNK_UPLOAD - cout << "loading ethash_hash_chunks" << endl; - m_hash_kernel = cl::Kernel(program, "ethash_hash_chunks"); - cout << "loading ethash_search_chunks" << endl; - m_search_kernel = cl::Kernel(program, "ethash_search_chunks"); - #else - cout << "loading ethash_hash" << endl; - m_hash_kernel = cl::Kernel(program, "ethash_hash"); - cout << "loading ethash_search" << endl; - m_search_kernel = cl::Kernel(program, "ethash_search"); - #endif - - // create buffer for dag - #ifdef ETHASH_CL_CHUNK_UPLOAD - for (unsigned i = 0; i < 4; i++) + if (_dagChunksNum == 1) + { + cout << "loading ethash_hash" << endl; + m_hash_kernel = cl::Kernel(program, "ethash_hash"); + cout << "loading ethash_search" << endl; + m_search_kernel = cl::Kernel(program, "ethash_search"); + } + else { - cout << "Creating chunky buffer: " << i << endl; - m_dags[i] = cl::Buffer(m_context, CL_MEM_READ_ONLY, (i == 3) ? (_dagSize - 3 * ((_dagSize >> 9) << 7)) : (_dagSize >> 9) << 7); + cout << "loading ethash_hash_chunks" << endl; + m_hash_kernel = cl::Kernel(program, "ethash_hash_chunks"); + cout << "loading ethash_search_chunks" << endl; + m_search_kernel = cl::Kernel(program, "ethash_search_chunks"); } - #else + + // create buffer for dag + if (_dagChunksNum == 1) + m_dagChunks[0] = cl::Buffer(m_context, CL_MEM_READ_ONLY, _dagSize); + else + for (unsigned i = 0; i < _dagChunksNum; i++) + { + // TODO Note: If we ever change to _dagChunksNum other than 4, then the size would need recalculation + cout << "Creating buffer for chunk " << i << endl; + m_dagChunks[i] = cl::Buffer( + m_context, + CL_MEM_READ_ONLY, + (i == 3) ? (_dagSize - 3 * ((_dagSize >> 9) << 7)) : (_dagSize >> 9) << 7 + ); + } cout << "Creating one big buffer." << endl; - m_dag = cl::Buffer(m_context, CL_MEM_READ_ONLY, _dagSize); - #endif // create buffer for header cout << "Creating buffer for header." << endl; m_header = cl::Buffer(m_context, CL_MEM_READ_ONLY, 32); - #ifdef ETHASH_CL_CHUNK_UPLOAD - void* dag_ptr[4]; - for (unsigned i = 0; i < 4; i++) + if (_dagChunksNum == 1) { - cout << "Mapping chunk " << i << endl; - dag_ptr[i] = m_queue.enqueueMapBuffer(m_dags[i], true, m_opencl_1_1 ? CL_MAP_WRITE : CL_MAP_WRITE_INVALIDATE_REGION, 0, (i == 3) ? (_dagSize - 3 * ((_dagSize >> 9) << 7)) : (_dagSize >> 9) << 7); + cout << "Mapping one big chunk." << endl; + m_queue.enqueueWriteBuffer(m_dagChunks[0], CL_TRUE, 0, _dagSize, _dag); } - for (unsigned i = 0; i < 4; i++) + else { - memcpy(dag_ptr[i], (char *)_dag + i*((_dagSize >> 9) << 7), (i == 3) ? (_dagSize - 3 * ((_dagSize >> 9) << 7)) : (_dagSize >> 9) << 7); - m_queue.enqueueUnmapMemObject(m_dags[i], dag_ptr[i]); + // TODO Note: If we ever change to _dagChunksNum other than 4, then the size would need recalculation + void* dag_ptr[4]; + for (unsigned i = 0; i < _dagChunksNum; i++) + { + cout << "Mapping chunk " << i << endl; + dag_ptr[i] = m_queue.enqueueMapBuffer(m_dagChunks[i], true, m_opencl_1_1 ? CL_MAP_WRITE : CL_MAP_WRITE_INVALIDATE_REGION, 0, (i == 3) ? (_dagSize - 3 * ((_dagSize >> 9) << 7)) : (_dagSize >> 9) << 7); + } + for (unsigned i = 0; i < _dagChunksNum; i++) + { + memcpy(dag_ptr[i], (char *)_dag + i*((_dagSize >> 9) << 7), (i == 3) ? (_dagSize - 3 * ((_dagSize >> 9) << 7)) : (_dagSize >> 9) << 7); + m_queue.enqueueUnmapMemObject(m_dagChunks[i], dag_ptr[i]); + } } - #else - cout << "Mapping chunk." << endl; - m_queue.enqueueWriteBuffer(m_dag, CL_TRUE, 0, _dagSize, _dag); - #endif // create mining buffers for (unsigned i = 0; i != c_num_buffers; ++i) @@ -288,22 +312,13 @@ void ethash_cl_miner::hash(uint8_t* ret, uint8_t const* header, uint64_t nonce, uint isolate ) */ - #ifdef ETHASH_CL_CHUNK_UPLOAD cout << "Setting chunk hash arguments." << endl; + unsigned argPos = 2; m_hash_kernel.setArg(1, m_header); - m_hash_kernel.setArg(2, m_dags[0]); - m_hash_kernel.setArg(3, m_dags[1]); - m_hash_kernel.setArg(4, m_dags[2]); - m_hash_kernel.setArg(5, m_dags[3]); - m_hash_kernel.setArg(6, nonce); - m_hash_kernel.setArg(7, ~0u); // have to pass this to stop the compile unrolling the loop - #else - cout << "Setting hash arguments." << endl; - m_hash_kernel.setArg(1, m_header); - m_hash_kernel.setArg(2, m_dag); - m_hash_kernel.setArg(3, nonce); - m_hash_kernel.setArg(4, ~0u); // have to pass this to stop the compile unrolling the loop - #endif + for (unsigned i = 0 ; i < m_dagChunksNum; ++i, ++argPos) + m_hash_kernel.setArg(argPos, m_dagChunks[i]); + m_hash_kernel.setArg(argPos + 1, nonce); + m_hash_kernel.setArg(argPos + 2, ~0u); // have to pass this to stop the compiler unrolling the loop unsigned buf = 0; for (unsigned i = 0; i < count || !pending.empty(); ) @@ -381,27 +396,13 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook uint isolate // 5 ) */ - #ifdef ETHASH_CL_CHUNK_UPLOAD - cout << "Setting chunk search arguments." << endl; - m_search_kernel.setArg(1, m_header); - m_search_kernel.setArg(2, m_dags[0]); - m_search_kernel.setArg(3, m_dags[1]); - m_search_kernel.setArg(4, m_dags[2]); - m_search_kernel.setArg(5, m_dags[3]); - - // pass these to stop the compiler unrolling the loops - m_search_kernel.setArg(7, target); - m_search_kernel.setArg(8, ~0u); - - #else - cout << "Setting search arguments." << endl; + unsigned argPos = 2; m_search_kernel.setArg(1, m_header); - m_search_kernel.setArg(2, m_dag); - + for (unsigned i = 0; i < m_dagChunksNum; ++i, ++argPos) + m_search_kernel.setArg(argPos, m_dagChunks[i]); // pass these to stop the compiler unrolling the loops - m_search_kernel.setArg(4, target); - m_search_kernel.setArg(5, ~0u); - #endif + m_search_kernel.setArg(argPos + 1, target); + m_search_kernel.setArg(argPos + 2, ~0u); unsigned buf = 0; std::random_device engine; @@ -410,11 +411,10 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook { // supply output buffer to kernel m_search_kernel.setArg(0, m_search_buf[buf]); - #ifdef ETHASH_CL_CHUNK_UPLOAD - m_search_kernel.setArg(6, start_nonce); - #else - m_search_kernel.setArg(3, start_nonce); - #endif + if (m_dagChunksNum == 1) + m_search_kernel.setArg(3, start_nonce); + else + m_search_kernel.setArg(6, start_nonce); // execute it! m_queue.enqueueNDRangeKernel(m_search_kernel, cl::NullRange, c_search_batch_size, m_workgroup_size); diff --git a/libethash-cl/ethash_cl_miner.h b/libethash-cl/ethash_cl_miner.h index 9c97f2aa4..0f83f8565 100644 --- a/libethash-cl/ethash_cl_miner.h +++ b/libethash-cl/ethash_cl_miner.h @@ -36,7 +36,14 @@ public: static unsigned get_num_devices(unsigned _platformId = 0); static std::string platform_info(unsigned _platformId = 0, unsigned _deviceId = 0); - bool init(uint8_t const* _dag, uint64_t _dagSize, unsigned workgroup_size = 64, unsigned _platformId = 0, unsigned _deviceId = 0); + bool init( + uint8_t const* _dag, + uint64_t _dagSize, + unsigned workgroup_size = 64, + unsigned _platformId = 0, + unsigned _deviceId = 0, + unsigned _dagChunksNum = 1 + ); void finish(); void hash(uint8_t* ret, uint8_t const* header, uint64_t nonce, unsigned count); void search(uint8_t const* header, uint64_t target, search_hook& hook); @@ -51,11 +58,11 @@ private: cl::CommandQueue m_queue; cl::Kernel m_hash_kernel; cl::Kernel m_search_kernel; - cl::Buffer m_dag; - cl::Buffer m_dags[4]; + unsigned m_dagChunksNum; + cl::Buffer* m_dagChunks; cl::Buffer m_header; cl::Buffer m_hash_buf[c_num_buffers]; cl::Buffer m_search_buf[c_num_buffers]; unsigned m_workgroup_size; bool m_opencl_1_1; -}; \ No newline at end of file +}; diff --git a/libethcore/Ethash.cpp b/libethcore/Ethash.cpp index f62c1f9cd..158f40981 100644 --- a/libethcore/Ethash.cpp +++ b/libethcore/Ethash.cpp @@ -285,6 +285,7 @@ private: unsigned Ethash::GPUMiner::s_platformId = 0; unsigned Ethash::GPUMiner::s_deviceId = 0; unsigned Ethash::GPUMiner::s_numInstances = 0; +unsigned Ethash::GPUMiner::s_dagChunks = 1; Ethash::GPUMiner::GPUMiner(ConstructionInfo const& _ci): Miner(_ci), @@ -345,7 +346,7 @@ void Ethash::GPUMiner::workLoop() this_thread::sleep_for(chrono::milliseconds(500)); } bytesConstRef dagData = dag->data(); - m_miner->init(dagData.data(), dagData.size(), 32, s_platformId, device); + m_miner->init(dagData.data(), dagData.size(), 32, s_platformId, device, s_dagChunks); } uint64_t upper64OfBoundary = (uint64_t)(u64)((u256)w.boundary >> 192); diff --git a/libethcore/Ethash.h b/libethcore/Ethash.h index 86540678f..868c27916 100644 --- a/libethcore/Ethash.h +++ b/libethcore/Ethash.h @@ -119,6 +119,7 @@ public: static void setDefaultPlatform(unsigned _id) { s_platformId = _id; } static void setDefaultDevice(unsigned _id) { s_deviceId = _id; } static void setNumInstances(unsigned _instances) { s_numInstances = std::min(_instances, getNumDevices()); } + static void setDagChunks(unsigned _dagChunks) { s_dagChunks = _dagChunks; } protected: void kickOff() override; @@ -137,6 +138,7 @@ public: static unsigned s_platformId; static unsigned s_deviceId; static unsigned s_numInstances; + static unsigned s_dagChunks; }; #else using GPUMiner = CPUMiner; From fdb06f08508df1b9358222d0785f3da41e014a92 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pawe=C5=82=20Bylica?= Date: Wed, 3 Jun 2015 17:17:41 +0200 Subject: [PATCH 09/20] Style correction. --- libethereum/Precompiled.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/libethereum/Precompiled.cpp b/libethereum/Precompiled.cpp index ab0d4d83f..20f89c452 100644 --- a/libethereum/Precompiled.cpp +++ b/libethereum/Precompiled.cpp @@ -83,12 +83,12 @@ void identityCode(bytesConstRef _in, bytesRef _out) std::unordered_map const& dev::eth::precompiled() { - static const std::unordered_map s_precompiled = + static const std::unordered_map c_precompiled = { { 1, { [](bytesConstRef) -> bigint { return c_ecrecoverGas; }, ecrecoverCode }}, { 2, { [](bytesConstRef i) -> bigint { return c_sha256Gas + (i.size() + 31) / 32 * c_sha256WordGas; }, sha256Code }}, { 3, { [](bytesConstRef i) -> bigint { return c_ripemd160Gas + (i.size() + 31) / 32 * c_ripemd160WordGas; }, ripemd160Code }}, { 4, { [](bytesConstRef i) -> bigint { return c_identityGas + (i.size() + 31) / 32 * c_identityWordGas; }, identityCode }} }; - return s_precompiled; + return c_precompiled; } From 65ebc5f17b061aa5976a82d3c42c970fa988034a Mon Sep 17 00:00:00 2001 From: Lefteris Karapetsas Date: Thu, 4 Jun 2015 12:22:27 +0200 Subject: [PATCH 10/20] Tweaks after merging latest develop changes - Use a vector for dag chunks. - Use ETHCL_LOG for outputing to stdout. --- libethash-cl/ethash_cl_miner.cpp | 61 ++++++++++---------------------- libethash-cl/ethash_cl_miner.h | 2 +- 2 files changed, 20 insertions(+), 43 deletions(-) diff --git a/libethash-cl/ethash_cl_miner.cpp b/libethash-cl/ethash_cl_miner.cpp index b2b7c49c2..aa90cb4c3 100644 --- a/libethash-cl/ethash_cl_miner.cpp +++ b/libethash-cl/ethash_cl_miner.cpp @@ -62,7 +62,7 @@ static void add_definition(std::string& source, char const* id, unsigned value) ethash_cl_miner::search_hook::~search_hook() {} ethash_cl_miner::ethash_cl_miner() -: m_dagChunks(nullptr), m_opencl_1_1() +: m_opencl_1_1() { } @@ -172,9 +172,6 @@ void ethash_cl_miner::finish() { if (m_queue()) m_queue.finish(); - - if (m_dagChunks) - delete [] m_dagChunks; } bool ethash_cl_miner::init( @@ -189,7 +186,6 @@ bool ethash_cl_miner::init( // for now due to the .cl kernels we can only have either 1 big chunk or 4 chunks assert(_dagChunksNum == 1 || _dagChunksNum == 4); // now create the number of chunk buffers - m_dagChunks = new cl::Buffer[_dagChunksNum]; m_dagChunksNum = _dagChunksNum; // get all platforms @@ -254,52 +250,52 @@ bool ethash_cl_miner::init( try { program.build({ device }); - cout << "Printing program log" << endl; - cout << program.getBuildInfo(device).c_str(); + ETHCL_LOG("Printing program log"); + ETHCL_LOG(program.getBuildInfo(device).c_str()); } catch (cl::Error err) { - cout << program.getBuildInfo(device).c_str(); + ETHCL_LOG(program.getBuildInfo(device).c_str()); return false; } if (_dagChunksNum == 1) { - cout << "loading ethash_hash" << endl; + ETHCL_LOG("Loading single big chunk kernels"); m_hash_kernel = cl::Kernel(program, "ethash_hash"); - cout << "loading ethash_search" << endl; m_search_kernel = cl::Kernel(program, "ethash_search"); } else { - cout << "loading ethash_hash_chunks" << endl; + ETHCL_LOG("Loading chunk kernels"); m_hash_kernel = cl::Kernel(program, "ethash_hash_chunks"); - cout << "loading ethash_search_chunks" << endl; m_search_kernel = cl::Kernel(program, "ethash_search_chunks"); } // create buffer for dag if (_dagChunksNum == 1) - m_dagChunks[0] = cl::Buffer(m_context, CL_MEM_READ_ONLY, _dagSize); + { + ETHCL_LOG("Creating one big buffer"); + m_dagChunks.push_back(cl::Buffer(m_context, CL_MEM_READ_ONLY, _dagSize)); + } else for (unsigned i = 0; i < _dagChunksNum; i++) { // TODO Note: If we ever change to _dagChunksNum other than 4, then the size would need recalculation - cout << "Creating buffer for chunk " << i << endl; - m_dagChunks[i] = cl::Buffer( + ETHCL_LOG("Creating buffer for chunk " << i); + m_dagChunks.push_back(cl::Buffer( m_context, CL_MEM_READ_ONLY, (i == 3) ? (_dagSize - 3 * ((_dagSize >> 9) << 7)) : (_dagSize >> 9) << 7 - ); + )); } - cout << "Creating one big buffer." << endl; // create buffer for header - cout << "Creating buffer for header." << endl; + ETHCL_LOG("Creating buffer for header."); m_header = cl::Buffer(m_context, CL_MEM_READ_ONLY, 32); if (_dagChunksNum == 1) { - cout << "Mapping one big chunk." << endl; + ETHCL_LOG("Mapping one big chunk."); m_queue.enqueueWriteBuffer(m_dagChunks[0], CL_TRUE, 0, _dagSize, _dag); } else @@ -308,7 +304,7 @@ bool ethash_cl_miner::init( void* dag_ptr[4]; for (unsigned i = 0; i < _dagChunksNum; i++) { - cout << "Mapping chunk " << i << endl; + ETHCL_LOG("Mapping chunk " << i); dag_ptr[i] = m_queue.enqueueMapBuffer(m_dagChunks[i], true, m_opencl_1_1 ? CL_MAP_WRITE : CL_MAP_WRITE_INVALIDATE_REGION, 0, (i == 3) ? (_dagSize - 3 * ((_dagSize >> 9) << 7)) : (_dagSize >> 9) << 7); } for (unsigned i = 0; i < _dagChunksNum; i++) @@ -321,7 +317,7 @@ bool ethash_cl_miner::init( // create mining buffers for (unsigned i = 0; i != c_num_buffers; ++i) { - cout << "Creating mining buffer " << i < m_dagChunks; cl::Buffer m_header; cl::Buffer m_hash_buf[c_num_buffers]; cl::Buffer m_search_buf[c_num_buffers]; From ee43c4f8a2ec35866a2615066f102467564330f1 Mon Sep 17 00:00:00 2001 From: chriseth Date: Sat, 6 Jun 2015 15:32:15 +0200 Subject: [PATCH 11/20] Fix: Optimiser was enabled by default. --- solc/CommandLineInterface.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/solc/CommandLineInterface.cpp b/solc/CommandLineInterface.cpp index 860b4c3b8..e65c602ab 100644 --- a/solc/CommandLineInterface.cpp +++ b/solc/CommandLineInterface.cpp @@ -412,8 +412,6 @@ bool CommandLineInterface::processInput() // TODO: Perhaps we should not compile unless requested bool optimize = m_args["optimize"].as(); unsigned runs = m_args["optimize-runs"].as(); - if (m_args.count("optimize-runs")) - optimize = true; m_compiler->compile(optimize, runs); } catch (ParserError const& _exception) From 9db5fb5bb6ff978123368a4406b71f1c5e249ab8 Mon Sep 17 00:00:00 2001 From: chriseth Date: Sat, 6 Jun 2015 01:04:55 +0200 Subject: [PATCH 12/20] Improved "Stack too deep" error message. Closes #2080. --- libevmasm/CommonSubexpressionEliminator.cpp | 4 ++-- libsolidity/ArrayUtils.cpp | 5 +++- libsolidity/Compiler.cpp | 2 +- libsolidity/CompilerUtils.cpp | 11 +++++---- libsolidity/ExpressionCompiler.cpp | 2 +- libsolidity/LValue.cpp | 26 +++++++++++++++------ 6 files changed, 34 insertions(+), 16 deletions(-) diff --git a/libevmasm/CommonSubexpressionEliminator.cpp b/libevmasm/CommonSubexpressionEliminator.cpp index b2fa73116..a441bd8bb 100644 --- a/libevmasm/CommonSubexpressionEliminator.cpp +++ b/libevmasm/CommonSubexpressionEliminator.cpp @@ -428,7 +428,7 @@ void CSECodeGenerator::appendDup(int _fromPosition, SourceLocation const& _locat { assertThrow(_fromPosition != c_invalidPosition, OptimizerException, ""); int instructionNum = 1 + m_stackHeight - _fromPosition; - assertThrow(instructionNum <= 16, StackTooDeepException, "Stack too deep."); + assertThrow(instructionNum <= 16, StackTooDeepException, "Stack too deep, try removing local variables."); assertThrow(1 <= instructionNum, OptimizerException, "Invalid stack access."); appendItem(AssemblyItem(dupInstruction(instructionNum), _location)); m_stack[m_stackHeight] = m_stack[_fromPosition]; @@ -441,7 +441,7 @@ void CSECodeGenerator::appendOrRemoveSwap(int _fromPosition, SourceLocation cons if (_fromPosition == m_stackHeight) return; int instructionNum = m_stackHeight - _fromPosition; - assertThrow(instructionNum <= 16, StackTooDeepException, "Stack too deep."); + assertThrow(instructionNum <= 16, StackTooDeepException, "Stack too deep, try removing local variables."); assertThrow(1 <= instructionNum, OptimizerException, "Invalid stack access."); appendItem(AssemblyItem(swapInstruction(instructionNum), _location)); diff --git a/libsolidity/ArrayUtils.cpp b/libsolidity/ArrayUtils.cpp index 531ab8af1..b770b815c 100644 --- a/libsolidity/ArrayUtils.cpp +++ b/libsolidity/ArrayUtils.cpp @@ -168,7 +168,10 @@ void ArrayUtils::copyArrayToStorage(ArrayType const& _targetType, ArrayType cons else solAssert(false, "Copying of unknown type requested: " + sourceBaseType->toString()); // stack: target_ref target_data_end source_data_pos target_data_pos source_data_end [target_byte_offset] [source_byte_offset] ... - solAssert(2 + byteOffsetSize + sourceBaseType->getSizeOnStack() <= 16, "Stack too deep."); + solAssert( + 2 + byteOffsetSize + sourceBaseType->getSizeOnStack() <= 16, + "Stack too deep, try removing local variables." + ); // fetch target storage reference m_context << eth::dupInstruction(2 + byteOffsetSize + sourceBaseType->getSizeOnStack()); if (haveByteOffsetTarget) diff --git a/libsolidity/Compiler.cpp b/libsolidity/Compiler.cpp index 40ed1fd8e..0a75e55a9 100644 --- a/libsolidity/Compiler.cpp +++ b/libsolidity/Compiler.cpp @@ -367,7 +367,7 @@ bool Compiler::visit(FunctionDefinition const& _function) stackLayout.push_back(i); stackLayout += vector(c_localVariablesSize, -1); - solAssert(stackLayout.size() <= 17, "Stack too deep."); + solAssert(stackLayout.size() <= 17, "Stack too deep, try removing local variables."); while (stackLayout.back() != int(stackLayout.size() - 1)) if (stackLayout.back() < 0) { diff --git a/libsolidity/CompilerUtils.cpp b/libsolidity/CompilerUtils.cpp index e3d8f7f90..3549ef98d 100644 --- a/libsolidity/CompilerUtils.cpp +++ b/libsolidity/CompilerUtils.cpp @@ -142,22 +142,25 @@ void CompilerUtils::moveToStackVariable(VariableDeclaration const& _variable) solAssert(stackPosition >= size, "Variable size and position mismatch."); // move variable starting from its top end in the stack if (stackPosition - size + 1 > 16) - BOOST_THROW_EXCEPTION(CompilerError() << errinfo_sourceLocation(_variable.getLocation()) - << errinfo_comment("Stack too deep.")); + BOOST_THROW_EXCEPTION( + CompilerError() << + errinfo_sourceLocation(_variable.getLocation()) << + errinfo_comment("Stack too deep, try removing local variables.") + ); for (unsigned i = 0; i < size; ++i) m_context << eth::swapInstruction(stackPosition - size + 1) << eth::Instruction::POP; } void CompilerUtils::copyToStackTop(unsigned _stackDepth, unsigned _itemSize) { - solAssert(_stackDepth <= 16, "Stack too deep."); + solAssert(_stackDepth <= 16, "Stack too deep, try removing local variables."); for (unsigned i = 0; i < _itemSize; ++i) m_context << eth::dupInstruction(_stackDepth); } void CompilerUtils::moveToStackTop(unsigned _stackDepth) { - solAssert(_stackDepth <= 15, "Stack too deep."); + solAssert(_stackDepth <= 15, "Stack too deep, try removing local variables."); for (unsigned i = 0; i < _stackDepth; ++i) m_context << eth::swapInstruction(1 + i); } diff --git a/libsolidity/ExpressionCompiler.cpp b/libsolidity/ExpressionCompiler.cpp index bac967d84..ba80a8ea2 100644 --- a/libsolidity/ExpressionCompiler.cpp +++ b/libsolidity/ExpressionCompiler.cpp @@ -263,7 +263,7 @@ bool ExpressionCompiler::visit(Assignment const& _assignment) appendOrdinaryBinaryOperatorCode(Token::AssignmentToBinaryOp(op), *_assignment.getType()); if (lvalueSize > 0) { - solAssert(itemSize + lvalueSize <= 16, "Stack too deep."); + solAssert(itemSize + lvalueSize <= 16, "Stack too deep, try removing local variables."); // value [lvalue_ref] updated_value for (unsigned i = 0; i < itemSize; ++i) m_context << eth::swapInstruction(itemSize + lvalueSize) << eth::Instruction::POP; diff --git a/libsolidity/LValue.cpp b/libsolidity/LValue.cpp index 38efb2a73..b684e55a1 100644 --- a/libsolidity/LValue.cpp +++ b/libsolidity/LValue.cpp @@ -42,8 +42,11 @@ void StackVariable::retrieveValue(SourceLocation const& _location, bool) const { unsigned stackPos = m_context.baseToCurrentStackOffset(m_baseStackOffset); if (stackPos >= 15) //@todo correct this by fetching earlier or moving to memory - BOOST_THROW_EXCEPTION(CompilerError() - << errinfo_sourceLocation(_location) << errinfo_comment("Stack too deep.")); + BOOST_THROW_EXCEPTION( + CompilerError() << + errinfo_sourceLocation(_location) << + errinfo_comment("Stack too deep, try removing local variables.") + ); for (unsigned i = 0; i < m_size; ++i) m_context << eth::dupInstruction(stackPos + 1); } @@ -52,8 +55,11 @@ void StackVariable::storeValue(Type const&, SourceLocation const& _location, boo { unsigned stackDiff = m_context.baseToCurrentStackOffset(m_baseStackOffset) - m_size + 1; if (stackDiff > 16) - BOOST_THROW_EXCEPTION(CompilerError() - << errinfo_sourceLocation(_location) << errinfo_comment("Stack too deep.")); + BOOST_THROW_EXCEPTION( + CompilerError() << + errinfo_sourceLocation(_location) << + errinfo_comment("Stack too deep, try removing local variables.") + ); else if (stackDiff > 0) for (unsigned i = 0; i < m_size; ++i) m_context << eth::swapInstruction(stackDiff) << eth::Instruction::POP; @@ -65,8 +71,11 @@ void StackVariable::setToZero(SourceLocation const& _location, bool) const { unsigned stackDiff = m_context.baseToCurrentStackOffset(m_baseStackOffset); if (stackDiff > 16) - BOOST_THROW_EXCEPTION(CompilerError() - << errinfo_sourceLocation(_location) << errinfo_comment("Stack too deep.")); + BOOST_THROW_EXCEPTION( + CompilerError() << + errinfo_sourceLocation(_location) << + errinfo_comment("Stack too deep, try removing local variables.") + ); solAssert(stackDiff >= m_size - 1, ""); for (unsigned i = 0; i < m_size; ++i) m_context << u256(0) << eth::swapInstruction(stackDiff + 1 - i) @@ -204,7 +213,10 @@ void StorageItem::storeValue(Type const& _sourceType, SourceLocation const& _loc // stack: source_ref source_off target_ref target_off member_slot_offset member_byte_offset source_member_ref source_member_off StorageItem(m_context, *memberType).retrieveValue(_location, true); // stack: source_ref source_off target_ref target_off member_offset source_value... - solAssert(4 + memberType->getSizeOnStack() <= 16, "Stack too deep."); + solAssert( + 4 + memberType->getSizeOnStack() <= 16, + "Stack too deep, try removing local varibales." + ); m_context << eth::dupInstruction(4 + memberType->getSizeOnStack()) << eth::dupInstruction(3 + memberType->getSizeOnStack()) << eth::Instruction::ADD From df8ca0d1a1ac749f76d34885e0c60a5c797f6611 Mon Sep 17 00:00:00 2001 From: Lefteris Karapetsas Date: Mon, 8 Jun 2015 00:56:13 +0200 Subject: [PATCH 13/20] Reclaim chunksNum set in MinerAux.h after merge --- ethminer/MinerAux.h | 1 + 1 file changed, 1 insertion(+) diff --git a/ethminer/MinerAux.h b/ethminer/MinerAux.h index 515dc2389..6a42dd774 100644 --- a/ethminer/MinerAux.h +++ b/ethminer/MinerAux.h @@ -268,6 +268,7 @@ public: ProofOfWork::GPUMiner::setDefaultPlatform(openclPlatform); ProofOfWork::GPUMiner::setDefaultDevice(openclDevice); ProofOfWork::GPUMiner::setNumInstances(miningThreads); + ProofOfWork::GPUMiner::setDagChunks(dagChunks); } if (mode == OperationMode::DAGInit) doInitDAG(initDAG); From c6020f6625fb98c0e96dd4f1946a7cd6b06877e9 Mon Sep 17 00:00:00 2001 From: Lefteris Karapetsas Date: Mon, 8 Jun 2015 01:22:17 +0200 Subject: [PATCH 14/20] Don't read kernel file as string during runtime --- libethash-cl/ethash_cl_miner.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/libethash-cl/ethash_cl_miner.cpp b/libethash-cl/ethash_cl_miner.cpp index 3d20ecc1b..f501d9642 100644 --- a/libethash-cl/ethash_cl_miner.cpp +++ b/libethash-cl/ethash_cl_miner.cpp @@ -233,9 +233,9 @@ bool ethash_cl_miner::init( m_workgroup_size = ((workgroup_size + 7) / 8) * 8; // patch source code - std::ifstream t("ethash_cl_miner_kernel.cl"); - std::string code((std::istreambuf_iterator(t)), - std::istreambuf_iterator()); + // note: ETHASH_CL_MINER_KERNEL is simply ethash_cl_miner_kernel.cl compiled + // into a byte array by bin2h.cmake. There is no need to load the file by hand in runtime + std::string code(ETHASH_CL_MINER_KERNEL, ETHASH_CL_MINER_KERNEL + ETHASH_CL_MINER_KERNEL_SIZE); add_definition(code, "GROUP_SIZE", m_workgroup_size); add_definition(code, "DAG_SIZE", (unsigned)(_dagSize / ETHASH_MIX_BYTES)); add_definition(code, "ACCESSES", ETHASH_ACCESSES); From f07ce95f78a369e2122d63d0f94b747f48b887fc Mon Sep 17 00:00:00 2001 From: Lefteris Karapetsas Date: Mon, 8 Jun 2015 01:33:17 +0200 Subject: [PATCH 15/20] setDagChunks is a part of the Interface --- libethcore/Ethash.h | 1 + 1 file changed, 1 insertion(+) diff --git a/libethcore/Ethash.h b/libethcore/Ethash.h index 48bf765d6..68c21c609 100644 --- a/libethcore/Ethash.h +++ b/libethcore/Ethash.h @@ -89,6 +89,7 @@ public: static std::string platformInfo(); static bool haveSufficientGPUMemory() { return false; } static void setDefaultPlatform(unsigned) {} + static void setDagChunks(unsigned) {} static void setDefaultDevice(unsigned) {} static void setNumInstances(unsigned _instances) { s_numInstances = std::min(_instances, std::thread::hardware_concurrency()); } protected: From 3a550d6ac777d6961466e5471faa39dffa0dfa69 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pawe=C5=82=20Bylica?= Date: Mon, 8 Jun 2015 10:27:37 +0200 Subject: [PATCH 16/20] Fix input data validation for ECRECOVER precompiled contract. --- libdevcrypto/Common.cpp | 2 +- libdevcrypto/Common.h | 2 +- libethereum/Precompiled.cpp | 5 +++-- 3 files changed, 5 insertions(+), 4 deletions(-) diff --git a/libdevcrypto/Common.cpp b/libdevcrypto/Common.cpp index 814f8309e..e68381427 100644 --- a/libdevcrypto/Common.cpp +++ b/libdevcrypto/Common.cpp @@ -37,7 +37,7 @@ using namespace dev::crypto; static Secp256k1 s_secp256k1; -bool dev::SignatureStruct::isValid() const +bool dev::SignatureStruct::isValid() const noexcept { if (v > 1 || r >= h256("0xfffffffffffffffffffffffffffffffebaaedce6af48a03bbfd25e8cd0364141") || diff --git a/libdevcrypto/Common.h b/libdevcrypto/Common.h index 10bcdd067..7bb51e563 100644 --- a/libdevcrypto/Common.h +++ b/libdevcrypto/Common.h @@ -51,7 +51,7 @@ struct SignatureStruct operator Signature() const { return *(h520 const*)this; } /// @returns true if r,s,v values are valid, otherwise false - bool isValid() const; + bool isValid() const noexcept; h256 r; h256 s; diff --git a/libethereum/Precompiled.cpp b/libethereum/Precompiled.cpp index 20f89c452..0ae9bf3eb 100644 --- a/libethereum/Precompiled.cpp +++ b/libethereum/Precompiled.cpp @@ -47,9 +47,10 @@ void ecrecoverCode(bytesConstRef _in, bytesRef _out) memcpy(&in, _in.data(), min(_in.size(), sizeof(in))); h256 ret; - if ((u256)in.v <= 28) + u256 v = (u256)in.v; + if (v >= 27 && v <= 28) { - SignatureStruct sig(in.r, in.s, (byte)((int)(u256)in.v - 27)); + SignatureStruct sig(in.r, in.s, (byte)((int)v - 27)); if (sig.isValid()) { try From f7ed7ed4defbf89758ed017c59ff2eb3c66f01f6 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pawe=C5=82=20Bylica?= Date: Mon, 8 Jun 2015 10:45:42 +0200 Subject: [PATCH 17/20] Rename Executive::collectResult -> setResultRecipient. --- libethereum/Client.cpp | 2 +- libethereum/Executive.h | 2 +- libethereum/State.cpp | 2 +- mix/MixClient.cpp | 2 +- test/libsolidity/solidityExecutionFramework.h | 2 +- 5 files changed, 5 insertions(+), 5 deletions(-) diff --git a/libethereum/Client.cpp b/libethereum/Client.cpp index 1e01a3a8c..f410e6eca 100644 --- a/libethereum/Client.cpp +++ b/libethereum/Client.cpp @@ -423,7 +423,7 @@ ExecutionResult Client::call(Address _dest, bytes const& _data, u256 _gas, u256 temp = m_postMine; temp.addBalance(_from, _value + _gasPrice * _gas); Executive e(temp, LastHashes(), 0); - e.collectResult(ret); + e.setResultRecipient(ret); if (!e.call(_dest, _from, _value, _gasPrice, &_data, _gas)) e.go(); e.finalize(); diff --git a/libethereum/Executive.h b/libethereum/Executive.h index 400c04298..beeee3331 100644 --- a/libethereum/Executive.h +++ b/libethereum/Executive.h @@ -119,7 +119,7 @@ public: bool excepted() const { return m_excepted != TransactionException::None; } /// Collect execution results in the result storage provided. - void collectResult(ExecutionResult& _res) { m_res = &_res; } + void setResultRecipient(ExecutionResult& _res) { m_res = &_res; } private: State& m_s; ///< The state to which this operation/transaction is applied. diff --git a/libethereum/State.cpp b/libethereum/State.cpp index 5d88405ac..20f080fdc 100644 --- a/libethereum/State.cpp +++ b/libethereum/State.cpp @@ -1190,7 +1190,7 @@ ExecutionResult State::execute(LastHashes const& _lh, Transaction const& _t, Per // transaction is bad in any way. Executive e(*this, _lh, 0); ExecutionResult res; - e.collectResult(res); + e.setResultRecipient(res); e.initialize(_t); // Uncommitting is a non-trivial operation - only do it once we've verified as much of the diff --git a/mix/MixClient.cpp b/mix/MixClient.cpp index 32875f572..7c0498ef0 100644 --- a/mix/MixClient.cpp +++ b/mix/MixClient.cpp @@ -132,7 +132,7 @@ void MixClient::executeTransaction(Transaction const& _t, State& _state, bool _c execState.addBalance(t.sender(), t.gas() * t.gasPrice()); //give it enough balance for gas estimation eth::ExecutionResult er; Executive execution(execState, lastHashes, 0); - execution.collectResult(er); + execution.setResultRecipient(er); execution.initialize(t); execution.execute(); std::vector machineStates; diff --git a/test/libsolidity/solidityExecutionFramework.h b/test/libsolidity/solidityExecutionFramework.h index 133b23604..a36e0b4ed 100644 --- a/test/libsolidity/solidityExecutionFramework.h +++ b/test/libsolidity/solidityExecutionFramework.h @@ -149,7 +149,7 @@ protected: m_state.addBalance(m_sender, _value); // just in case eth::Executive executive(m_state, eth::LastHashes(), 0); eth::ExecutionResult res; - executive.collectResult(res); + executive.setResultRecipient(res); eth::Transaction t = _isCreation ? eth::Transaction(_value, m_gasPrice, m_gas, _data, 0, KeyPair::create().sec()) : From e81fc1e68fd8fe233b1b0a88c64bf58d3677c9aa Mon Sep 17 00:00:00 2001 From: Gav Wood Date: Mon, 8 Jun 2015 23:28:46 +0900 Subject: [PATCH 18/20] Minor debug alterations/fixes for blockchain downloading to make it play marginally better with the braindead Go strategy. Import key without address. --- alethzero/MainWin.cpp | 8 +++--- ethkey/KeyAux.h | 39 ++++++++++++++++++++++++++- libdevcore/CommonData.cpp | 52 +++++++++++++++++------------------- libdevcore/CommonData.h | 2 +- libdevcrypto/SecretStore.cpp | 2 +- libethcore/KeyManager.cpp | 23 +++++++++------- libethcore/KeyManager.h | 11 +++++--- libethereum/CommonNet.h | 6 +++-- libethereum/EthereumHost.cpp | 4 ++- libethereum/EthereumPeer.cpp | 4 +-- 10 files changed, 98 insertions(+), 53 deletions(-) diff --git a/alethzero/MainWin.cpp b/alethzero/MainWin.cpp index 10038a80e..93054fd67 100644 --- a/alethzero/MainWin.cpp +++ b/alethzero/MainWin.cpp @@ -2114,14 +2114,14 @@ void Main::on_reencryptKey_triggered() auto pw = [&](){ auto p = QInputDialog::getText(this, "Re-Encrypt Key", "Enter the original password for this key.\nHint: " + QString::fromStdString(m_keyManager.hint(a)), QLineEdit::Password, QString()).toStdString(); if (p.empty()) - throw UnknownPassword(); + throw PasswordUnknown(); return p; }; while (!(password.empty() ? m_keyManager.recode(a, SemanticPassword::Master, pw, kdf) : m_keyManager.recode(a, password, hint, pw, kdf))) if (QMessageBox::question(this, "Re-Encrypt Key", "Password given is incorrect. Would you like to try again?", QMessageBox::Retry, QMessageBox::Cancel) == QMessageBox::Cancel) return; } - catch (UnknownPassword&) {} + catch (PasswordUnknown&) {} } } @@ -2137,13 +2137,13 @@ void Main::on_reencryptAll_triggered() while (!m_keyManager.recode(a, SemanticPassword::Existing, [&](){ auto p = QInputDialog::getText(nullptr, "Re-Encrypt Key", QString("Enter the original password for key %1.\nHint: %2").arg(QString::fromStdString(pretty(a))).arg(QString::fromStdString(m_keyManager.hint(a))), QLineEdit::Password, QString()).toStdString(); if (p.empty()) - throw UnknownPassword(); + throw PasswordUnknown(); return p; }, (KDF)kdfs.indexOf(kdf))) if (QMessageBox::question(this, "Re-Encrypt Key", "Password given is incorrect. Would you like to try again?", QMessageBox::Retry, QMessageBox::Cancel) == QMessageBox::Cancel) return; } - catch (UnknownPassword&) {} + catch (PasswordUnknown&) {} } void Main::on_go_triggered() diff --git a/ethkey/KeyAux.h b/ethkey/KeyAux.h index 639e1d4f4..d2ec13b2a 100644 --- a/ethkey/KeyAux.h +++ b/ethkey/KeyAux.h @@ -102,6 +102,7 @@ public: List, New, Import, + ImportWithAddress, Export, Recode, Kill @@ -159,6 +160,13 @@ public: m_inputs = strings(1, argv[++i]); m_name = argv[++i]; } + else if ((arg == "-i" || arg == "--import-with-address") && i + 3 < argc) + { + m_mode = OperationMode::ImportWithAddress; + m_inputs = strings(1, argv[++i]); + m_address = Address(argv[++i]); + m_name = argv[++i]; + } else if (arg == "--export") m_mode = OperationMode::Export; else if (arg == "--recode") @@ -314,6 +322,33 @@ public: cout << " ICAP: " << ICAP(k.address()).encoded() << endl; break; } + case OperationMode::ImportWithAddress: + { + string const& i = m_inputs[0]; + h128 u; + bytes b; + b = fromHex(i); + if (b.size() != 32) + { + std::string s = contentsString(i); + b = fromHex(s); + if (b.size() != 32) + u = wallet.store().importKey(i); + } + if (!u && b.size() == 32) + u = wallet.store().importSecret(b, lockPassword(toAddress(Secret(b)).abridged())); + if (!u) + { + cerr << "Cannot import " << i << " not a file or secret." << endl; + break; + } + wallet.importExisting(u, m_name, m_address); + cout << "Successfully imported " << i << ":" << endl; + cout << " Name: " << m_name << endl; + cout << " Address: " << m_address << endl; + cout << " UUID: " << toUUID(u) << endl; + break; + } case OperationMode::List: { vector bare; @@ -369,6 +404,7 @@ public: << " -l,--list List all keys available in wallet." << endl << " -n,--new Create a new key with given name and add it in the wallet." << endl << " -i,--import [||] Import keys from given source and place in wallet." << endl + << " --import-with-address [||]
Import keys from given source with given address and place in wallet." << endl << " -e,--export [
| , ... ] Export given keys." << endl << " -r,--recode [
|| , ... ] Decrypt and re-encrypt given keys." << endl << "Wallet configuration:" << endl @@ -418,8 +454,9 @@ private: string m_lockHint; bool m_icap = true; - /// Creating + /// Creating/importing string m_name; + Address m_address; /// Importing strings m_inputs; diff --git a/libdevcore/CommonData.cpp b/libdevcore/CommonData.cpp index f8d8c172f..2d6333f26 100644 --- a/libdevcore/CommonData.cpp +++ b/libdevcore/CommonData.cpp @@ -67,7 +67,7 @@ std::string dev::randomWord() return ret; } -int dev::fromHex(char _i) +int dev::fromHex(char _i, WhenError _throw) { if (_i >= '0' && _i <= '9') return _i - '0'; @@ -75,7 +75,10 @@ int dev::fromHex(char _i) return _i - 'a' + 10; if (_i >= 'A' && _i <= 'F') return _i - 'A' + 10; - BOOST_THROW_EXCEPTION(BadHexCharacter() << errinfo_invalidSymbol(_i)); + if (_throw == WhenError::Throw) + BOOST_THROW_EXCEPTION(BadHexCharacter() << errinfo_invalidSymbol(_i)); + else + return -1; } bytes dev::fromHex(std::string const& _s, WhenError _throw) @@ -85,33 +88,26 @@ bytes dev::fromHex(std::string const& _s, WhenError _throw) ret.reserve((_s.size() - s + 1) / 2); if (_s.size() % 2) - try - { - ret.push_back(fromHex(_s[s++])); - } - catch (...) - { - ret.push_back(0); - // msvc does not support it -#ifndef BOOST_NO_EXCEPTIONS - cwarn << boost::current_exception_diagnostic_information(); -#endif - if (_throw == WhenError::Throw) - throw; - } + { + int h = fromHex(_s[s++], WhenError::DontThrow); + if (h != -1) + ret.push_back(h); + else if (_throw == WhenError::Throw) + throw BadHexCharacter(); + else + return bytes(); + } for (unsigned i = s; i < _s.size(); i += 2) - try - { - ret.push_back((byte)(fromHex(_s[i]) * 16 + fromHex(_s[i + 1]))); - } - catch (...){ - ret.push_back(0); -#ifndef BOOST_NO_EXCEPTIONS - cwarn << boost::current_exception_diagnostic_information(); -#endif - if (_throw == WhenError::Throw) - throw; - } + { + int h = fromHex(_s[i], WhenError::DontThrow); + int l = fromHex(_s[i + 1], WhenError::DontThrow); + if (h != -1 && l != -1) + ret.push_back((byte)(h * 16 + l)); + else if (_throw == WhenError::Throw) + throw BadHexCharacter(); + else + return bytes(); + } return ret; } diff --git a/libdevcore/CommonData.h b/libdevcore/CommonData.h index e1d8d7bdb..ddc00e09f 100644 --- a/libdevcore/CommonData.h +++ b/libdevcore/CommonData.h @@ -61,7 +61,7 @@ std::string toHex(_T const& _data, int _w = 2, HexPrefix _prefix = HexPrefix::Do /// Converts a (printable) ASCII hex character into the correspnding integer value. /// @example fromHex('A') == 10 && fromHex('f') == 15 && fromHex('5') == 5 -int fromHex(char _i); +int fromHex(char _i, WhenError _throw); /// Converts a (printable) ASCII hex string into the corresponding byte stream. /// @example fromHex("41626261") == asBytes("Abba") diff --git a/libdevcrypto/SecretStore.cpp b/libdevcrypto/SecretStore.cpp index 11ff98bf6..b9d4ccfc6 100644 --- a/libdevcrypto/SecretStore.cpp +++ b/libdevcrypto/SecretStore.cpp @@ -164,7 +164,7 @@ void SecretStore::load(std::string const& _keysPath) h128 SecretStore::readKey(std::string const& _file, bool _deleteFile) { - cdebug << "Reading" << _file; + cnote << "Reading" << _file; return readKeyContent(contentsString(_file), _deleteFile ? _file : string()); } diff --git a/libethcore/KeyManager.cpp b/libethcore/KeyManager.cpp index 182201301..4430a588e 100644 --- a/libethcore/KeyManager.cpp +++ b/libethcore/KeyManager.cpp @@ -89,18 +89,18 @@ bool KeyManager::load(std::string const& _pass) for (auto const& i: s[1]) { m_keyInfo[m_addrLookup[(Address)i[0]] = (h128)i[1]] = KeyInfo((h256)i[2], (std::string)i[3]); - cdebug << toString((Address)i[0]) << toString((h128)i[1]) << toString((h256)i[2]) << (std::string)i[3]; +// cdebug << toString((Address)i[0]) << toString((h128)i[1]) << toString((h256)i[2]) << (std::string)i[3]; } for (auto const& i: s[2]) m_passwordInfo[(h256)i[0]] = (std::string)i[1]; m_password = (string)s[3]; } - cdebug << hashPassword(m_password) << toHex(m_password); +// cdebug << hashPassword(m_password) << toHex(m_password); m_cachedPasswords[hashPassword(m_password)] = m_password; - cdebug << hashPassword(asString(m_key.ref())) << m_key.hex(); +// cdebug << hashPassword(asString(m_key.ref())) << m_key.hex(); m_cachedPasswords[hashPassword(asString(m_key.ref()))] = asString(m_key.ref()); - cdebug << hashPassword(_pass) << _pass; +// cdebug << hashPassword(_pass) << _pass; m_cachedPasswords[m_master = hashPassword(_pass)] = _pass; return true; } @@ -141,7 +141,7 @@ std::string KeyManager::getPassword(h256 const& _passHash, function const& _pass = DontKnowThrow) const; Secret secret(h128 const& _uuid, std::function const& _pass = DontKnowThrow) const; diff --git a/libethereum/CommonNet.h b/libethereum/CommonNet.h index a2f4a2e7c..8cf2647cf 100644 --- a/libethereum/CommonNet.h +++ b/libethereum/CommonNet.h @@ -38,14 +38,16 @@ namespace eth #if ETH_DEBUG static const unsigned c_maxHashes = 2048; ///< Maximum number of hashes BlockHashes will ever send. -static const unsigned c_maxHashesAsk = 2048; ///< Maximum number of hashes GetBlockHashes will ever ask for. +static const unsigned c_maxHashesAsk = 256; ///< Maximum number of hashes GetBlockHashes will ever ask for. static const unsigned c_maxBlocks = 128; ///< Maximum number of blocks Blocks will ever send. -static const unsigned c_maxBlocksAsk = 128; ///< Maximum number of blocks we ask to receive in Blocks (when using GetChain). +static const unsigned c_maxBlocksAsk = 8; ///< Maximum number of blocks we ask to receive in Blocks (when using GetChain). +static const unsigned c_maxPayload = 262144; ///< Maximum size of packet for us to send. #else static const unsigned c_maxHashes = 2048; ///< Maximum number of hashes BlockHashes will ever send. static const unsigned c_maxHashesAsk = 2048; ///< Maximum number of hashes GetBlockHashes will ever ask for. static const unsigned c_maxBlocks = 128; ///< Maximum number of blocks Blocks will ever send. static const unsigned c_maxBlocksAsk = 128; ///< Maximum number of blocks we ask to receive in Blocks (when using GetChain). +static const unsigned c_maxPayload = 262144; ///< Maximum size of packet for us to send. #endif class BlockChain; diff --git a/libethereum/EthereumHost.cpp b/libethereum/EthereumHost.cpp index 3965a214e..7c9b730ea 100644 --- a/libethereum/EthereumHost.cpp +++ b/libethereum/EthereumHost.cpp @@ -253,7 +253,9 @@ void EthereumHost::onPeerStatus(EthereumPeer* _peer) if (_peer->m_protocolVersion != protocolVersion()) estimatePeerHashes(_peer); else if (_peer->m_latestBlockNumber > m_chain.number()) - _peer->m_expectedHashes = (unsigned)_peer->m_latestBlockNumber - m_chain.number(); + _peer->m_expectedHashes = (unsigned)_peer->m_latestBlockNumber - m_chain.number() + 1000; + else + _peer->m_expectedHashes = 1000; if (m_hashMan.chainSize() < _peer->m_expectedHashes) m_hashMan.resetToRange(m_chain.number() + 1, _peer->m_expectedHashes); continueSync(_peer); diff --git a/libethereum/EthereumPeer.cpp b/libethereum/EthereumPeer.cpp index d6b0b50c3..a332e5b93 100644 --- a/libethereum/EthereumPeer.cpp +++ b/libethereum/EthereumPeer.cpp @@ -263,7 +263,7 @@ bool EthereumPeer::interpret(unsigned _id, RLP const& _r) // return the requested blocks. bytes rlp; unsigned n = 0; - for (unsigned i = 0; i < min(count, c_maxBlocks); ++i) + for (unsigned i = 0; i < min(count, c_maxBlocks) && rlp.size() < c_maxPayload; ++i) { auto h = _r[i].toHash(); if (host()->chain().isKnown(h)) @@ -286,7 +286,7 @@ bool EthereumPeer::interpret(unsigned _id, RLP const& _r) case BlocksPacket: { if (m_asking != Asking::Blocks) - clog(NetWarn) << "Peer giving us blocks when we didn't ask for them."; + clog(NetImpolite) << "Peer giving us blocks when we didn't ask for them."; else { setAsking(Asking::Nothing); From 1f9330a29c3791d4a4ae5e6e0561e4e6f8cd97c1 Mon Sep 17 00:00:00 2001 From: Gav Wood Date: Mon, 8 Jun 2015 23:54:16 +0900 Subject: [PATCH 19/20] Avoid divide by zero. --- alethzero/Transact.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/alethzero/Transact.cpp b/alethzero/Transact.cpp index e61092f33..fd466e475 100644 --- a/alethzero/Transact.cpp +++ b/alethzero/Transact.cpp @@ -360,7 +360,7 @@ void Transact::rejigData() return; } else - gasNeeded = (qint64)min(ethereum()->gasLimitRemaining(), ((b - value()) / gasPrice())); + gasNeeded = (qint64)min(ethereum()->gasLimitRemaining(), ((b - value()) / max(gasPrice(), 1))); // Dry-run execution to determine gas requirement and any execution errors Address to; From 3589c73b75937a4cba202e872d7b0269f035448a Mon Sep 17 00:00:00 2001 From: Gav Wood Date: Mon, 8 Jun 2015 23:54:42 +0900 Subject: [PATCH 20/20] Version bump. --- libdevcore/Common.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/libdevcore/Common.cpp b/libdevcore/Common.cpp index cb9b94de8..3dc3fd280 100644 --- a/libdevcore/Common.cpp +++ b/libdevcore/Common.cpp @@ -28,7 +28,7 @@ using namespace dev; namespace dev { -char const* Version = "0.9.23"; +char const* Version = "0.9.24"; const u256 UndefinedU256 = ~(u256)0;