ethashCL doesn't need memory to shadow each GPU.

libethash-cl/ethash_cl_miner.cpp

@@ -88,6 +88,13 @@ std::string ethash_cl_miner::platform_info(unsigned _platformId, unsigned _devic
 	return "{ \"platform\": \"" + platforms[platform_num].getInfo<CL_PLATFORM_NAME>() + "\", \"device\": \"" + device.getInfo<CL_DEVICE_NAME>() + "\", \"version\": \"" + device_version + "\" }";
 }
 
+unsigned ethash_cl_miner::get_num_platforms()
+{
+	std::vector<cl::Platform> platforms;
+	cl::Platform::get(&platforms);
+	return platforms.size();
+}
+
 unsigned ethash_cl_miner::get_num_devices(unsigned _platformId)
 {
 	std::vector<cl::Platform> platforms;
@@ -117,14 +124,11 @@ void ethash_cl_miner::finish()
 	}
 }
 
-bool ethash_cl_miner::init(ethash_params const& params, std::function<void(void*)> _fillDAG, 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)
 {
-	// store params
-	m_params = params;
-
 	// get all platforms
-    std::vector<cl::Platform> platforms;
+	std::vector<cl::Platform> platforms;
-    cl::Platform::get(&platforms);
+	cl::Platform::get(&platforms);
 	if (platforms.empty())
 	{
 		cout << "No OpenCL platforms found." << endl;
@@ -137,10 +141,10 @@ bool ethash_cl_miner::init(ethash_params const& params, std::function<void(void*
 
 	cout << "Using platform: " << platforms[_platformId].getInfo<CL_PLATFORM_NAME>().c_str() << endl;
 
-    // get GPU device of the default platform
+	// get GPU device of the default platform
-    std::vector<cl::Device> devices;
+	std::vector<cl::Device> devices;
 	platforms[_platformId].getDevices(CL_DEVICE_TYPE_ALL, &devices);
-    if (devices.empty())
+	if (devices.empty())
 	{
 		cout << "No OpenCL devices found." << endl;
 		return false;
@@ -171,7 +175,7 @@ bool ethash_cl_miner::init(ethash_params const& params, std::function<void(void*
 	// 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)(params.full_size / ETHASH_MIX_BYTES));
+	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());
@@ -194,18 +198,20 @@ bool ethash_cl_miner::init(ethash_params const& params, std::function<void(void*
 	m_search_kernel = cl::Kernel(program, "ethash_search");
 
 	// create buffer for dag
-	m_dag = cl::Buffer(m_context, CL_MEM_READ_ONLY, params.full_size);
+	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
 	{
+		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, params.full_size);
+//		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);
+//		_fillDAG(dag_ptr);
-		m_queue.enqueueUnmapMemObject(m_dag, dag_ptr);
+//		m_queue.enqueueUnmapMemObject(m_dag, dag_ptr);
 	}
 
 	// create mining buffers

libethash-cl/ethash_cl_miner.h

@@ -31,11 +31,11 @@ public:
 public:
 	ethash_cl_miner();
 
-	bool init(ethash_params const& params, std::function<void(void*)> _fillDAG, unsigned workgroup_size = 64, unsigned _platformId = 0, unsigned _deviceId = 0);
+	static unsigned get_num_platforms();
-	static std::string platform_info(unsigned _platformId = 0, unsigned _deviceId = 0);
 	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);
 	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);
@@ -43,7 +43,6 @@ public:
 private:
 	enum { c_max_search_results = 63, c_num_buffers = 2, c_hash_batch_size = 1024, c_search_batch_size = 1024*256 };
 
-	ethash_params m_params;
 	cl::Context m_context;
 	cl::CommandQueue m_queue;
 	cl::Kernel m_hash_kernel;

libethcore/Ethash.cpp

@@ -309,10 +309,10 @@ void Ethash::GPUMiner::workLoop()
 			delete m_miner;
 			m_miner = new ethash_cl_miner;
 
-			auto p = EthashAux::params(m_minerSeed);
-			auto cb = [&](void* d) { EthashAux::full(m_minerSeed, bytesRef((byte*)d, p.full_size)); };
 			unsigned device = instances() > 1 ? index() : s_deviceId;
-			m_miner->init(p, cb, 32, s_platformId, device);
+
+			EthashAux::FullType dag = EthashAux::full(m_minerSeed);
+			m_miner->init(dag->data.data(), dag->data.size(), 32, s_platformId, device);
 		}
 
 		uint64_t upper64OfBoundary = (uint64_t)(u64)((u256)w.boundary >> 192);

libethcore/EthashAux.cpp

@@ -133,7 +133,6 @@ EthashAux::LightAllocation::~LightAllocation()
 	ethash_delete_light(light);
 }
 
-
 EthashAux::FullType EthashAux::full(BlockInfo const& _header, bytesRef _dest, bool _createIfMissing)
 {
 	return full(_header.seedHash(), _dest, _createIfMissing);