Remove --use-chunks argument and general refactoring

- Removing --use-chunks argument. The decision to use chunks or not is
  now made by the implementation, depending on
  CL_DEVICE_MAX_MEM_ALLOC_SIZE

- Refactored the code a bit in ethash_cl_miner, abstracting out some of
  the device iteration into its own functions.

ethminer/MinerAux.h

@@ -132,10 +132,6 @@ public:
 			ProofOfWork::GPUMiner::listDevices();
 			exit(0);
 		}
-		else if (arg == "--use-chunks")
-		{
-			dagChunks = 4;
-		}
 		else if (arg == "--phone-home" && i + 1 < argc)
 		{
 			string m = argv[++i];
@@ -180,7 +176,7 @@ public:
 			m_minerType = MinerType::CPU;
 		else if (arg == "-G" || arg == "--opencl")
 		{
-			if (!ProofOfWork::GPUMiner::haveSufficientMemory())
+			if (!ProofOfWork::GPUMiner::configureGPU())
 			{
 				cout << "No GPU device with sufficient memory was found. Defaulting to CPU" << endl;
 				m_minerType = MinerType::CPU;
@@ -273,7 +269,6 @@ public:
 			ProofOfWork::GPUMiner::setDefaultPlatform(openclPlatform);
 			ProofOfWork::GPUMiner::setDefaultDevice(openclDevice);
 			ProofOfWork::GPUMiner::setNumInstances(miningThreads);
-			ProofOfWork::GPUMiner::setDagChunks(dagChunks);
 		}
 		if (mode == OperationMode::DAGInit)
 			doInitDAG(initDAG);
@@ -311,7 +306,6 @@ public:
 			<< "    --opencl-platform <n>  When mining using -G/--opencl use OpenCL platform n (default: 0)." << endl
 			<< "    --opencl-device <n>  When mining using -G/--opencl use OpenCL device n (default: 0)." << endl
 			<< "    -t, --mining-threads <n> 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
 			;
 	}
 

libethash-cl/ethash_cl_miner.cpp

@@ -127,7 +127,7 @@ unsigned ethash_cl_miner::getNumDevices(unsigned _platformId)
 	return devices.size();
 }
 
-bool ethash_cl_miner::haveSufficientGPUMemory()
+bool ethash_cl_miner::configureGPU()
 {
 	return searchForAllDevices([](cl::Device const _device) -> bool
 		{
@@ -184,15 +184,40 @@ bool ethash_cl_miner::searchForAllDevices(unsigned _platformId, std::function<bo
 	return false;
 }
 
+void ethash_cl_miner::doForAllDevices(std::function<void(cl::Device const&)> _callback)
+{
+	std::vector<cl::Platform> platforms;
+	cl::Platform::get(&platforms);
+	if (platforms.empty())
+	{
+		ETHCL_LOG("No OpenCL platforms found.");
+		return;
+	}
+	for (unsigned i = 0; i < platforms.size(); ++i)
+		doForAllDevices(i, _callback);
+}
+
+void ethash_cl_miner::doForAllDevices(unsigned _platformId, std::function<void(cl::Device const&)> _callback)
+{
+	std::vector<cl::Platform> platforms;
+	cl::Platform::get(&platforms);
+	if (_platformId >= platforms.size())
+		return;
+
+	std::vector<cl::Device> devices;
+	platforms[_platformId].getDevices(CL_DEVICE_TYPE_ALL, &devices);
+	for (cl::Device const& device: devices)
+		_callback(device);
+}
+
 void ethash_cl_miner::listDevices()
 {
 	std::string outString ="\nListing OpenCL devices.\nFORMAT: [deviceID] deviceName\n";
 	unsigned int i = 0;
-	searchForAllDevices([&outString, &i](cl::Device const _device) -> bool
+	doForAllDevices([&outString, &i](cl::Device const _device)
 		{
 			outString += "[" + to_string(i) + "] " + _device.getInfo<CL_DEVICE_NAME>() + "\n";
 			++i;
-			return false; // so that search continues
 		}
 	);
 	ETHCL_LOG(outString);
@@ -209,15 +234,9 @@ bool ethash_cl_miner::init(
 	uint64_t _dagSize,
 	unsigned workgroup_size,
 	unsigned _platformId,
-	unsigned _deviceId,
-	unsigned _dagChunksNum
+	unsigned _deviceId
 )
 {
-	// 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_dagChunksNum = _dagChunksNum;
-
 	// get all platforms
 	try
 	{
@@ -247,6 +266,10 @@ bool ethash_cl_miner::init(
 		std::string device_version = device.getInfo<CL_DEVICE_VERSION>();
 		ETHCL_LOG("Using device: " << device.getInfo<CL_DEVICE_NAME>().c_str() << "(" << device_version.c_str() << ")");
 
+		// configure chunk number depending on max allocateable memory
+		cl_ulong result;		
+		device.getInfo(CL_DEVICE_MAX_MEM_ALLOC_SIZE, &result);
+		m_dagChunksNum = result >= ETHASH_CL_MINIMUM_MEMORY ? 4 : 1;
 		if (strncmp("OpenCL 1.0", device_version.c_str(), 10) == 0)
 		{
 			ETHCL_LOG("OpenCL 1.0 is not supported.");
@@ -288,7 +311,7 @@ bool ethash_cl_miner::init(
 			ETHCL_LOG(program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(device).c_str());
 			return false;
 		}
-		if (_dagChunksNum == 1)
+		if (m_dagChunksNum == 1)
 		{
 			ETHCL_LOG("Loading single big chunk kernels");
 			m_hash_kernel = cl::Kernel(program, "ethash_hash");
@@ -302,13 +325,13 @@ bool ethash_cl_miner::init(
 		}
 
 		// create buffer for dag
-		if (_dagChunksNum == 1)
+		if (m_dagChunksNum == 1)
 		{
 			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++)
+			for (unsigned i = 0; i < m_dagChunksNum; i++)
 			{
 				// TODO Note: If we ever change to _dagChunksNum other than 4, then the size would need recalculation
 				ETHCL_LOG("Creating buffer for chunk " << i);
@@ -323,7 +346,7 @@ bool ethash_cl_miner::init(
 		ETHCL_LOG("Creating buffer for header.");
 		m_header = cl::Buffer(m_context, CL_MEM_READ_ONLY, 32);
 
-		if (_dagChunksNum == 1)
+		if (m_dagChunksNum == 1)
 		{
 			ETHCL_LOG("Mapping one big chunk.");
 			m_queue.enqueueWriteBuffer(m_dagChunks[0], CL_TRUE, 0, _dagSize, _dag);
@@ -332,12 +355,12 @@ bool ethash_cl_miner::init(
 		{
 			// 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++)
+			for (unsigned i = 0; i < m_dagChunksNum; i++)
 			{
 				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++)
+			for (unsigned i = 0; i < m_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]);

libethash-cl/ethash_cl_miner.h

@@ -34,10 +34,11 @@ public:
 
 	static bool searchForAllDevices(unsigned _platformId, std::function<bool(cl::Device const&)> _callback);
 	static bool searchForAllDevices(std::function<bool(cl::Device const&)> _callback);
+	static void doForAllDevices(unsigned _platformId, std::function<void(cl::Device const&)> _callback);
+	static void doForAllDevices(std::function<void(cl::Device const&)> _callback);
 	static unsigned getNumPlatforms();
 	static unsigned getNumDevices(unsigned _platformId = 0);
 	static std::string platform_info(unsigned _platformId = 0, unsigned _deviceId = 0);
-	static bool haveSufficientGPUMemory();
 	static void listDevices();
 	static bool configureGPU();
 
@@ -46,8 +47,7 @@ public:
 		uint64_t _dagSize,
 		unsigned workgroup_size = 64,
 		unsigned _platformId = 0,
-		unsigned _deviceId = 0,
-		unsigned _dagChunksNum = 1
+		unsigned _deviceId = 0
 	);
 	void finish();
 	void search(uint8_t const* header, uint64_t target, search_hook& hook);
@@ -62,7 +62,7 @@ private:
 	cl::CommandQueue m_queue;
 	cl::Kernel m_hash_kernel;
 	cl::Kernel m_search_kernel;
-	unsigned m_dagChunksNum;
+	unsigned int m_dagChunksNum;
 	std::vector<cl::Buffer> m_dagChunks;
 	cl::Buffer m_header;
 	cl::Buffer m_hash_buf[c_num_buffers];
@@ -70,5 +70,4 @@ private:
 	unsigned m_workgroup_size;
 	bool m_opencl_1_1;
 
-	static const unsigned int s_dagChunksNum = 1;
 };

libethcore/Ethash.cpp

@@ -285,7 +285,6 @@ 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),
@@ -346,7 +345,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, s_dagChunks);
+			m_miner->init(dagData.data(), dagData.size(), 32, s_platformId, device);
 		}
 
 		uint64_t upper64OfBoundary = (uint64_t)(u64)((u256)w.boundary >> 192);
@@ -379,9 +378,9 @@ void Ethash::GPUMiner::listDevices()
 	return ethash_cl_miner::listDevices();
 }
 
-bool Ethash::GPUMiner::haveSufficientMemory()
+bool Ethash::GPUMiner::configureGPU()
 {
-	return ethash_cl_miner::haveSufficientGPUMemory();
+	return ethash_cl_miner::configureGPU();
 }
 
 #endif

libethcore/Ethash.h

@@ -91,7 +91,7 @@ public:
 		static void setDagChunks(unsigned) {}
 		static void setDefaultDevice(unsigned) {}
 		static void listDevices() {}
-		static bool haveSufficientMemory() { return false; }
+		static bool configureGPU() { return false; }
 		static void setNumInstances(unsigned _instances) { s_numInstances = std::min<unsigned>(_instances, std::thread::hardware_concurrency()); }
 	protected:
 		void kickOff() override
@@ -120,11 +120,10 @@ public:
 		static std::string platformInfo();
 		static unsigned getNumDevices();
 		static void listDevices();
-		static bool haveSufficientMemory();
+		static bool configureGPU();
 		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<unsigned>(_instances, getNumDevices()); }
-		static void setDagChunks(unsigned _dagChunks) { s_dagChunks = _dagChunks; }
 
 	protected:
 		void kickOff() override;
@@ -143,7 +142,6 @@ public:
 		static unsigned s_platformId;
 		static unsigned s_deviceId;
 		static unsigned s_numInstances;
-		static unsigned s_dagChunks;
 	};
 #else
 	using GPUMiner = CPUMiner;