OpenCL: Always try single chunk DAG upload

- Removed the `--force-single-chunk` option

- Always attempt to create a single chunk DAG buffer in the GPU. If that
  fails then and only then switch to multiple chunks.

This change is motivated by the fact that many GPUs appear to be able to
actually allocate a lot more than what CL_DEVICE_MAX_MEM_ALLOC_SIZE
returns which proves that the results of querying the CL API on this
basically can't be trusted.

ethminer/MinerAux.h

@@ -134,8 +134,6 @@ 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];
@@ -273,7 +271,6 @@ public:
 					m_openclDevice,
 					m_clAllowCPU,
 					m_extraGPUMemory,
-					m_forceSingleChunk,
 					m_currentBlock
 				))
 			{
@@ -321,7 +318,6 @@ 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
 			;
 	}
 
@@ -510,7 +506,6 @@ 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

@@ -140,12 +140,10 @@ unsigned ethash_cl_miner::getNumDevices(unsigned _platformId)
 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;
@@ -174,7 +172,6 @@ bool ethash_cl_miner::configureGPU(
 }
 
 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)
@@ -288,23 +285,6 @@ bool ethash_cl_miner::init(
 		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);
-		if (s_forceSingleChunk || result >= _dagSize)
-		{
-			m_dagChunksNum = 1;
-			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);
-		}
-
 		if (strncmp("OpenCL 1.0", device_version.c_str(), 10) == 0)
 		{
 			ETHCL_LOG("OpenCL 1.0 is not supported.");
@@ -346,26 +326,23 @@ bool ethash_cl_miner::init(
 			ETHCL_LOG(program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(device).c_str());
 			return false;
 		}
-		if (m_dagChunksNum == 1)
-		{
-			ETHCL_LOG("Loading single big chunk kernels");
-			m_hash_kernel = cl::Kernel(program, "ethash_hash");
-			m_search_kernel = cl::Kernel(program, "ethash_search");
-		}
-		else
-		{
-			ETHCL_LOG("Loading chunk kernels");
-			m_hash_kernel = cl::Kernel(program, "ethash_hash_chunks");
-			m_search_kernel = cl::Kernel(program, "ethash_search_chunks");
-		}
 
 		// create buffer for dag
-		if (m_dagChunksNum == 1)
+		try
 		{
-			ETHCL_LOG("Creating one big buffer");
+			m_dagChunksNum = 1;
 			m_dagChunks.push_back(cl::Buffer(m_context, CL_MEM_READ_ONLY, _dagSize));
+			ETHCL_LOG("Created one big buffer for the DAG");
 		}
-		else
+		catch (...)
+		{
+			cl_ulong result;
+			device.getInfo(CL_DEVICE_MAX_MEM_ALLOC_SIZE, &result);
+			ETHCL_LOG(
+				"Failed to allocate 1 big chunk. Max allocateable memory is "
+				<< result << ". Trying to allocate 4 chunks."
+			);
+			m_dagChunksNum = 4;
 			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
@@ -376,6 +353,20 @@ bool ethash_cl_miner::init(
 					(i == 3) ? (_dagSize - 3 * ((_dagSize >> 9) << 7)) : (_dagSize >> 9) << 7
 				));
 			}
+		}
+
+		if (m_dagChunksNum == 1)
+		{
+			ETHCL_LOG("Loading single big chunk kernels");
+			m_hash_kernel = cl::Kernel(program, "ethash_hash");
+			m_search_kernel = cl::Kernel(program, "ethash_search");
+		}
+		else
+		{
+			ETHCL_LOG("Loading chunk kernels");
+			m_hash_kernel = cl::Kernel(program, "ethash_hash_chunks");
+			m_search_kernel = cl::Kernel(program, "ethash_search_chunks");
+		}
 
 		// create buffer for header
 		ETHCL_LOG("Creating buffer for header.");

libethash-cl/ethash_cl_miner.h

@@ -44,7 +44,6 @@ public:
 	static bool configureGPU(
 		bool _allowCPU,
 		unsigned _extraGPUMemory,
-		bool _forceSingleChunk,
 		boost::optional<uint64_t> _currentBlock
 	);
 
@@ -79,8 +78,6 @@ 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

@@ -389,13 +389,12 @@ 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, _forceSingleChunk, _currentBlock);
+	return ethash_cl_miner::configureGPU(_allowCPU, _extraGPUMemory, _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, bool, boost::optional<uint64_t>) { return false; }
+		static bool configureGPU(unsigned, unsigned, bool, unsigned,  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,7 +122,6 @@ 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()); }