Merge branch 'develop' of github.com:ethereum/cpp-ethereum into develop

cmake/EthCompilerSettings.cmake

@@ -55,11 +55,6 @@ else ()
 	message(WARNING "Your compiler is not tested, if you run into any issues, we'd welcome any patches.")
 endif ()
 
-# Set stack memory limits
-if (APPLE)
-	set(CMAKE_EXE_LINKER_FLAGS "${CMAKE_EXE_LINKER_FLAGS} -Wl,-stack_size,1000000")
-endif()
-
 if (PROFILING AND (("${CMAKE_CXX_COMPILER_ID}" MATCHES "GNU") OR ("${CMAKE_CXX_COMPILER_ID}" MATCHES "Clang")))
 	set(CMAKE_CXX_FLAGS "-g ${CMAKE_CXX_FLAGS}")
 	set(CMAKE_C_FLAGS "-g ${CMAKE_C_FLAGS}")

ethminer/MinerAux.h

@@ -134,6 +134,8 @@ public:
 			m_clAllowCPU = true;
 		else if (arg == "--cl-extragpu-mem" && i + 1 < argc)
 			m_extraGPUMemory = 1000000 * stol(argv[++i]);
+		else if (arg == "--force-single-chunk")
+			m_forceSingleChunk = true;
 		else if (arg == "--phone-home" && i + 1 < argc)
 		{
 			string m = argv[++i];
@@ -271,6 +273,7 @@ public:
 					m_openclDevice,
 					m_clAllowCPU,
 					m_extraGPUMemory,
+					m_forceSingleChunk,
 					m_currentBlock
 				))
 			{
@@ -318,6 +321,7 @@ public:
 			<< "    --list-devices List the detected OpenCL devices and exit." <<endl
 			<< "    --current-block Let the miner know the current block number at configuration time. Will help determine DAG size and required GPU memory." <<endl
 			<< "    --cl-extragpu-mem Set the memory (in MB) you believe your GPU requires for stuff other than mining. Windows rendering e.t.c.." <<endl
+			<< "    --force-single-chunk Force DAG uploading in a single chunk against OpenCL's judgement. Use at your own risk." <<endl
 			;
 	}
 
@@ -507,6 +511,7 @@ private:
 	unsigned m_miningThreads = UINT_MAX;
 	bool m_shouldListDevices = false;
 	bool m_clAllowCPU = false;
+	bool m_forceSingleChunk = false;
 	boost::optional<uint64_t> m_currentBlock;
 	// default value is 350MB of GPU memory for other stuff (windows system rendering, e.t.c.)
 	unsigned m_extraGPUMemory = 350000000;

libethash-cl/ethash_cl_miner.cpp

@@ -137,9 +137,15 @@ unsigned ethash_cl_miner::getNumDevices(unsigned _platformId)
 	return devices.size();
 }
 
-bool ethash_cl_miner::configureGPU(bool _allowCPU, unsigned _extraGPUMemory, boost::optional<uint64_t> _currentBlock)
+bool ethash_cl_miner::configureGPU(
+	bool _allowCPU,
+	unsigned _extraGPUMemory,
+	bool _forceSingleChunk,
+	boost::optional<uint64_t> _currentBlock
+)
 {
 	s_allowCPU = _allowCPU;
+	s_forceSingleChunk = _forceSingleChunk;
 	s_extraRequiredGPUMem = _extraGPUMemory;
 	// by default let's only consider the DAG of the first epoch
 	uint64_t dagSize = _currentBlock ? ethash_get_datasize(*_currentBlock) : 1073739904U;
@@ -168,6 +174,7 @@ bool ethash_cl_miner::configureGPU(bool _allowCPU, unsigned _extraGPUMemory, boo
 }
 
 bool ethash_cl_miner::s_allowCPU = false;
+bool ethash_cl_miner::s_forceSingleChunk = false;
 unsigned ethash_cl_miner::s_extraRequiredGPUMem;
 
 bool ethash_cl_miner::searchForAllDevices(function<bool(cl::Device const&)> _callback)
@@ -284,15 +291,18 @@ bool ethash_cl_miner::init(
 		// configure chunk number depending on max allocateable memory
 		cl_ulong result;
 		device.getInfo(CL_DEVICE_MAX_MEM_ALLOC_SIZE, &result);
-		if (result >= _dagSize)
+		if (s_forceSingleChunk || result >= _dagSize)
 		{
 			m_dagChunksNum = 1;
-			ETHCL_LOG("Using 1 big chunk. Max OpenCL allocateable memory is" << result);
+			ETHCL_LOG(
+				((result <= _dagSize && s_forceSingleChunk) ? "Forcing single chunk. Good luck!\n" : "") <<
+				"Using 1 big chunk. Max OpenCL allocateable memory is " << result
+			);
 		}
 		else
 		{
 			m_dagChunksNum = 4;
-			ETHCL_LOG("Using 4 chunks. Max OpenCL allocateable memory is" << result);
+			ETHCL_LOG("Using 4 chunks. Max OpenCL allocateable memory is " << result);
 		}
 
 		if (strncmp("OpenCL 1.0", device_version.c_str(), 10) == 0)

libethash-cl/ethash_cl_miner.h

@@ -41,7 +41,12 @@ public:
 	static unsigned getNumDevices(unsigned _platformId = 0);
 	static std::string platform_info(unsigned _platformId = 0, unsigned _deviceId = 0);
 	static void listDevices();
-	static bool configureGPU(bool _allowCPU, unsigned _extraGPUMemory, boost::optional<uint64_t> _currentBlock);
+	static bool configureGPU(
+		bool _allowCPU,
+		unsigned _extraGPUMemory,
+		bool _forceSingleChunk,
+		boost::optional<uint64_t> _currentBlock
+	);
 
 	bool init(
 		uint8_t const* _dag,
@@ -74,6 +79,8 @@ private:
 	unsigned m_workgroup_size;
 	bool m_opencl_1_1;
 
+	/// Force dag upload to GPU in a single chunk even if OpenCL thinks you can't do it. Use at your own risk.
+	static bool s_forceSingleChunk;
 	/// Allow CPU to appear as an OpenCL device or not. Default is false
 	static bool s_allowCPU;
 	/// GPU memory required for other things, like window rendering e.t.c.

libethcore/Ethash.cpp

@@ -386,12 +386,13 @@ bool Ethash::GPUMiner::configureGPU(
 	unsigned _deviceId,
 	bool _allowCPU,
 	unsigned _extraGPUMemory,
+	bool _forceSingleChunk,
 	boost::optional<uint64_t> _currentBlock
 )
 {
 	s_platformId = _platformId;
 	s_deviceId = _deviceId;
-	return ethash_cl_miner::configureGPU(_allowCPU, _extraGPUMemory, _currentBlock);
+	return ethash_cl_miner::configureGPU(_allowCPU, _extraGPUMemory, _forceSingleChunk, _currentBlock);
 }
 
 #endif

libethcore/Ethash.h

@@ -88,7 +88,7 @@ public:
 		static unsigned instances() { return s_numInstances > 0 ? s_numInstances : std::thread::hardware_concurrency(); }
 		static std::string platformInfo();
 		static void listDevices() {}
-		static bool configureGPU(unsigned, unsigned, bool, unsigned, boost::optional<uint64_t>) { return false; }
+		static bool configureGPU(unsigned, unsigned, bool, unsigned, bool, boost::optional<uint64_t>) { return false; }
 		static void setNumInstances(unsigned _instances) { s_numInstances = std::min<unsigned>(_instances, std::thread::hardware_concurrency()); }
 	protected:
 		void kickOff() override
@@ -122,6 +122,7 @@ public:
 			unsigned _deviceId,
 			bool _allowCPU,
 			unsigned _extraGPUMemory,
+			bool _forceSingleChunk,
 			boost::optional<uint64_t> _currentBlock
 		);
 		static void setNumInstances(unsigned _instances) { s_numInstances = std::min<unsigned>(_instances, getNumDevices()); }

libethereum/ExtVM.cpp

@@ -20,18 +20,69 @@
  */
 
 #include "ExtVM.h"
-
+#include <exception>
+#include <boost/thread.hpp>
 #include "Executive.h"
-using namespace std;
+
 using namespace dev;
 using namespace dev::eth;
 
+namespace
+{
+
+static unsigned const c_depthLimit = 1024;
+
+/// Upper bound of stack space needed by single CALL/CREATE execution. Set experimentally.
+static size_t const c_singleExecutionStackSize = 12 * 1024;
+
+/// Standard OSX thread stack limit. Should be reasonable for other platforms too.
+static size_t const c_defaultStackSize = 512 * 1024;
+
+/// On what depth execution should be offloaded to additional separated stack space.
+static unsigned const c_offloadPoint = c_defaultStackSize / c_singleExecutionStackSize;
+
+void goOnOffloadedStack(Executive& _e, OnOpFunc const& _onOp)
+{
+	// Set new stack size enouth to handle the rest of the calls up to the limit.
+	boost::thread::attributes attrs;
+	attrs.set_stack_size((c_depthLimit - c_offloadPoint) * c_singleExecutionStackSize);
+
+	// Create new thread with big stack and join immediately.
+	// TODO: It is possible to switch the implementation to Boost.Context or similar when the API is stable.
+	std::exception_ptr exception;
+	boost::thread{attrs, [&]{
+		try
+		{
+			_e.go(_onOp);
+		}
+		catch (...)
+		{
+			exception = std::current_exception(); // Catch all exceptions to be rethrown in parent thread.
+		}
+	}}.join();
+	if (exception)
+		std::rethrow_exception(exception);
+}
+
+void go(unsigned _depth, Executive& _e, OnOpFunc const& _onOp)
+{
+	// If in the offloading point we need to switch to additional separated stack space.
+	// Current stack is too small to handle more CALL/CREATE executions.
+	// It needs to be done only once as newly allocated stack space it enough to handle 
+	// the rest of the calls up to the depth limit (c_depthLimit).
+	if (_depth == c_offloadPoint)
+		goOnOffloadedStack(_e, _onOp);
+	else
+		_e.go(_onOp);
+}
+}
+
 bool ExtVM::call(CallParameters& _p)
 {
 	Executive e(m_s, lastHashes, depth + 1);
 	if (!e.call(_p, gasPrice, origin))
 	{
-		e.go(_p.onOp);
+		go(depth, e, _p.onOp);
 		e.accrueSubState(sub);
 	}
 	_p.gas = e.gas();
@@ -47,7 +98,7 @@ h160 ExtVM::create(u256 _endowment, u256& io_gas, bytesConstRef _code, OnOpFunc
 	Executive e(m_s, lastHashes, depth + 1);
 	if (!e.create(myAddress, _endowment, gasPrice, io_gas, _code, origin))
 	{
-		e.go(_onOp);
+		go(depth, e, _onOp);
 		e.accrueSubState(sub);
 	}
 	io_gas = e.gas();