diff --git a/CMakeLists.txt b/CMakeLists.txt
index 037728a9e..e7e89a5a7 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -131,7 +131,7 @@ configureProject()
 set (ETH_HAVE_WEBENGINE 1)
 
 message(STATUS "CMAKE_VERSION: ${CMAKE_VERSION}")
-message("-- VMTRACE: ${VMTRACE}; PARANOIA: ${PARANOIA}; HEADLESS: ${HEADLESS}; JSONRPC: ${JSONRPC}; EVMJIT: ${EVMJIT}; FATDB: ${FATDB}; CHROMIUM: ${ETH_HAVE_WEBENGINE}; USENPM: ${USENPM}")
+message("-- VMTRACE: ${VMTRACE}; PARANOIA: ${PARANOIA}; HEADLESS: ${HEADLESS}; JSONRPC: ${JSONRPC}; EVMJIT: ${EVMJIT}; FATDB: ${FATDB}; CHROMIUM: ${ETH_HAVE_WEBENGINE}; USENPM: ${USENPM}; ETHASHCL: ${ETHASHCL}")
 
 
 # Default TARGET_PLATFORM to "linux".
diff --git a/libethcore/BlockInfo.cpp b/libethcore/BlockInfo.cpp
index ef57264a1..f04353aac 100644
--- a/libethcore/BlockInfo.cpp
+++ b/libethcore/BlockInfo.cpp
@@ -61,7 +61,7 @@ void BlockInfo::setEmpty()
 	hash = headerHash(WithNonce);
 }
 
-h256 BlockInfo::seedHash() const
+h256 const& BlockInfo::seedHash() const
 {
 	if (!m_seedHash)
 		for (u256 n = number; n >= c_epochDuration; n -= c_epochDuration)
diff --git a/libethcore/BlockInfo.h b/libethcore/BlockInfo.h
index 912978ef0..8f42dc038 100644
--- a/libethcore/BlockInfo.h
+++ b/libethcore/BlockInfo.h
@@ -128,7 +128,7 @@ public:
 
 	u256 calculateDifficulty(BlockInfo const& _parent) const;
 	u256 selectGasLimit(BlockInfo const& _parent) const;
-	h256 seedHash() const;
+	h256 const& seedHash() const;
 
 	/// sha3 of the header only.
 	h256 headerHash(IncludeNonce _n) const;
diff --git a/libethcore/ProofOfWork.cpp b/libethcore/ProofOfWork.cpp
index d261ccb1c..3cbfb4905 100644
--- a/libethcore/ProofOfWork.cpp
+++ b/libethcore/ProofOfWork.cpp
@@ -30,6 +30,9 @@
 #include <libdevcrypto/CryptoPP.h>
 #include <libdevcrypto/FileSystem.h>
 #include <libdevcore/Common.h>
+#if ETH_ETHASHCL
+#include <libethash-cl/ethash_cl_miner.h>
+#endif
 #include "BlockInfo.h"
 #include "Ethasher.h"
 #include "ProofOfWork.h"
@@ -41,12 +44,12 @@ namespace dev
 namespace eth
 {
 
-bool Ethash::verify(BlockInfo const& _header)
+bool EthashCPU::verify(BlockInfo const& _header)
 {
 	return Ethasher::verify(_header);
 }
 
-std::pair<MineInfo, Ethash::Proof> Ethash::mine(BlockInfo const& _header, unsigned _msTimeout, bool _continue, bool _turbo)
+std::pair<MineInfo, EthashCPU::Proof> EthashCPU::mine(BlockInfo const& _header, unsigned _msTimeout, bool _continue, bool _turbo)
 {
 	Ethasher::Miner m(_header);
 
@@ -93,5 +96,113 @@ std::pair<MineInfo, Ethash::Proof> Ethash::mine(BlockInfo const& _header, unsign
 	return ret;
 }
 
+#if ETH_ETHASHCL
+
+/*
+struct ethash_cl_search_hook
+{
+	// reports progress, return true to abort
+	virtual bool found(uint64_t const* nonces, uint32_t count) = 0;
+	virtual bool searched(uint64_t start_nonce, uint32_t count) = 0;
+};
+
+class ethash_cl_miner
+{
+public:
+	ethash_cl_miner();
+
+	bool init(ethash_params const& params, const uint8_t seed[32], unsigned workgroup_size = 64);
+
+	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);
+};
+*/
+
+struct EthashCLHook: public ethash_cl_search_hook
+{
+	virtual bool found(uint64_t const* _nonces, uint32_t _count)
+	{
+		Guard l(x_all);
+		for (unsigned i = 0; i < _count; ++i)
+			found.push_back((Nonce)(u64)_nonces[i]);
+		if (abort)
+		{
+			aborted = true;
+			return true;
+		}
+		return false;
+	}
+
+	virtual bool searched(uint64_t _startNonce, uint32_t _count)
+	{
+		Guard l(x_all);
+		total += _count;
+		last = _startNonce + _count;
+		if (abort)
+		{
+			aborted = true;
+			return true;
+		}
+		return false;
+	}
+
+	vector<Nonce> fetchFound() { vector<Nonce> ret; Guard l(x_all); std::swap(ret, found); return ret; }
+	uint64_t fetchTotal() { Guard l(x_all); auto ret = total; total = 0; return ret; }
+
+	Mutex x_all;
+	vector<Nonce> found;
+	uint64_t total;
+	uint64_t last;
+	bool abort = false;
+	bool aborted = false;
+};
+
+EthashCL::EthashCL():
+	m_miner(new ethash_cl_miner),
+	m_hook(new EthashCLHook)
+{
+}
+
+EthashCL::~EthashCL()
+{
+	m_hook->abort = true;
+	for (unsigned timeout = 0; timeout < 100 && !m_hook->aborted; ++timeout)
+		std::this_thread::sleep_for(chrono::milliseconds(30));
+	if (!m_hook->aborted)
+		cwarn << "Couldn't abort. Abandoning OpenCL process.";
+}
+
+bool EthashCL::verify(BlockInfo const& _header)
+{
+	return Ethasher::verify(_header);
+}
+
+std::pair<MineInfo, Ethash::Proof> EthashCL::mine(BlockInfo const& _header, unsigned _msTimeout, bool, bool)
+{
+	if (m_lastHeader.seedHash() != _header.seedHash())
+	{
+		m_miner->init(Ethasher::params(_header), _header.seedHash().data());
+		// TODO: reinit probably won't work when seed changes.
+	}
+	if (m_lastHeader != _header)
+	{
+		static std::random_device s_eng;
+		uint64_t tryNonce = (uint64_t)(u64)(m_last = Nonce::random(s_eng));
+		m_miner->search(_header.headerHash(WithoutNonce).data(), tryNonce, *m_hook);
+	}
+	m_lastHeader = _header;
+
+	std::this_thread::sleep_for(chrono::milliseconds(_msTimeout));
+	auto found = m_hook->fetchFound();
+	if (!found.empty())
+	{
+		h256 mixHash; // ?????
+		return std::make_pair(MineInfo{0.0, 1e99, 0, true}, EthashCL::Proof((Nonce)(u64)found[0], mixHash));
+	}
+	return std::make_pair(MineInfo{0.0, 1e99, 0, false}, EthashCL::Proof());
+}
+
+#endif
+
 }
 }
diff --git a/libethcore/ProofOfWork.h b/libethcore/ProofOfWork.h
index 250ddb73d..245a96a02 100644
--- a/libethcore/ProofOfWork.h
+++ b/libethcore/ProofOfWork.h
@@ -32,6 +32,9 @@
 
 #define FAKE_DAGGER 1
 
+class ethash_cl_miner;
+struct ethash_cl_search_hook;
+
 namespace dev
 {
 namespace eth
@@ -46,7 +49,7 @@ struct MineInfo
 	bool completed = false;
 };
 
-class Ethash
+class EthashCPU
 {
 public:
 	struct Proof
@@ -63,6 +66,36 @@ protected:
 	Nonce m_last;
 };
 
+#if ETH_ETHASHCL
+class EthashCL
+{
+public:
+	struct Proof
+	{
+		Nonce nonce;
+		h256 mixHash;
+	};
+
+	EthashCL();
+	~EthashCL();
+
+	static bool verify(BlockInfo const& _header);
+	std::pair<MineInfo, Proof> mine(BlockInfo const& _header, unsigned _msTimeout = 100, bool _continue = true, bool _turbo = false);
+	static void assignResult(Proof const& _r, BlockInfo& _header) { _header.nonce = _r.nonce; _header.mixHash = _r.mixHash; }
+
+protected:
+	Nonce m_last;
+	BlockInfo m_lastHeader;
+	Nonce m_mined;
+	std::unique_ptr<ethash_cl_miner> m_miner;
+	std::unique_ptr<ethash_cl_search_hook> m_hook;
+};
+
+using Ethash = EthashCL;
+#else
+using Ethash = EthashCPU;
+#endif
+
 template <class Evaluator>
 class ProofOfWorkEngine: public Evaluator
 {
diff --git a/libethereum/CMakeLists.txt b/libethereum/CMakeLists.txt
index 94a4e4497..3df3b9bc3 100644
--- a/libethereum/CMakeLists.txt
+++ b/libethereum/CMakeLists.txt
@@ -36,6 +36,10 @@ target_link_libraries(${EXECUTABLE} devcrypto)
 target_link_libraries(${EXECUTABLE} ethcore)
 target_link_libraries(${EXECUTABLE} secp256k1)
 
+if (ETHASHCL)
+	target_link_libraries(${EXECUTABLE} ethash-cl)
+endif ()
+
 if (CMAKE_COMPILER_IS_MINGW)
 	target_link_libraries(${EXECUTABLE} ssp shlwapi)
 endif()
diff --git a/libethereum/Client.cpp b/libethereum/Client.cpp
index bd8b3da40..f46d0d3c9 100644
--- a/libethereum/Client.cpp
+++ b/libethereum/Client.cpp
@@ -411,8 +411,11 @@ void Client::setForceMining(bool _enable)
 void Client::setMiningThreads(unsigned _threads)
 {
 	stopMining();
-
+#if ETH_ETHASHCL
+	unsigned t = 1;
+#else
 	auto t = _threads ? _threads : thread::hardware_concurrency();
+#endif
 	WriteGuard l(x_localMiners);
 	m_localMiners.clear();
 	m_localMiners.resize(t);
diff --git a/libethereum/Client.h b/libethereum/Client.h
index cdfdcd1eb..04798ecdb 100644
--- a/libethereum/Client.h
+++ b/libethereum/Client.h
@@ -323,12 +323,12 @@ public:
 	virtual unsigned miningThreads() const { ReadGuard l(x_localMiners); return m_localMiners.size(); }
 	/// Start mining.
 	/// NOT thread-safe - call it & stopMining only from a single thread
-	virtual void startMining() { startWorking(); ReadGuard l(x_localMiners); for (auto& m: m_localMiners) m.start(); }
+	virtual void startMining() { startWorking(); { ReadGuard l(x_localMiners); for (auto& m: m_localMiners) m.start(); } }
 	/// Stop mining.
 	/// NOT thread-safe
-	virtual void stopMining() { ReadGuard l(x_localMiners); for (auto& m: m_localMiners) m.stop(); }
+	virtual void stopMining() { { ReadGuard l(x_localMiners); for (auto& m: m_localMiners) m.stop(); } }
 	/// Are we mining now?
-	virtual bool isMining() { ReadGuard l(x_localMiners); return m_localMiners.size() && m_localMiners[0].isRunning(); }
+	virtual bool isMining() { { ReadGuard l(x_localMiners); if (!m_localMiners.empty() && m_localMiners[0].isRunning()) return true; } return false; }
 	/// Check the progress of the mining.
 	virtual MineProgress miningProgress() const;
 	/// Get and clear the mining history.
@@ -405,8 +405,6 @@ private:
 	bool m_forceMining = false;				///< Mine even when there are no transactions pending?
 	bool m_verifyOwnBlocks = true;			///< Should be verify blocks that we mined?
 
-
-
 	mutable Mutex m_filterLock;
 	std::map<h256, InstalledFilter> m_filters;
 	std::map<unsigned, ClientWatch> m_watches;
diff --git a/libethereum/Miner.cpp b/libethereum/Miner.cpp
index a049fca2f..08bc426dc 100644
--- a/libethereum/Miner.cpp
+++ b/libethereum/Miner.cpp
@@ -31,15 +31,15 @@ using namespace dev::eth;
 Miner::~Miner() {}
 
 LocalMiner::LocalMiner(MinerHost* _host, unsigned _id):
-	Worker("miner-" + toString(_id)),
-	m_host(_host)
+	AsyncMiner(_host, _id),
+	Worker("miner-" + toString(_id))
 {
 }
 
 void LocalMiner::setup(MinerHost* _host, unsigned _id)
 {
-	m_host = _host;
-	setName("miner-" + toString(_id));
+	AsyncMiner::setup(_host, _id);
+	setName("miner-" + toString(m_id));
 }
 
 void LocalMiner::doWork()
diff --git a/libethereum/Miner.h b/libethereum/Miner.h
index e472c6f64..eaa9c9e8c 100644
--- a/libethereum/Miner.h
+++ b/libethereum/Miner.h
@@ -74,6 +74,32 @@ public:
 	virtual bytes const& blockData() const = 0;
 };
 
+class AsyncMiner: public Miner
+{
+public:
+	/// Null constructor.
+	AsyncMiner(): m_host(nullptr) {}
+
+	/// Constructor.
+	AsyncMiner(MinerHost* _host, unsigned _id = 0): m_host(_host), m_id(_id) {}
+
+	/// Setup its basics.
+	void setup(MinerHost* _host, unsigned _id = 0) { m_host = _host; m_id = _id; }
+
+	/// Start mining.
+	virtual void start() {}
+
+	/// Stop mining.
+	virtual void stop() {}
+
+	/// @returns true iff the mining has been start()ed. It may still not be actually mining, depending on the host's turbo() & force().
+	virtual bool isRunning() { return false; }
+
+protected:
+	MinerHost* m_host = nullptr;			///< Our host.
+	unsigned m_id = 0;						///< Our unique id.
+};
+
 /**
  * @brief Implements Miner.
  * To begin mining, use start() & stop(). noteStateChange() can be used to reset the mining and set up the
@@ -86,11 +112,11 @@ public:
  * @threadsafe
  * @todo Signal Miner to restart once with condition variables.
  */
-class LocalMiner: public Miner, Worker
+class LocalMiner: public AsyncMiner, Worker
 {
 public:
 	/// Null constructor.
-	LocalMiner(): m_host(nullptr) {}
+	LocalMiner() {}
 
 	/// Constructor.
 	LocalMiner(MinerHost* _host, unsigned _id = 0);