removed --cuda-turbo option, implemented a new --cuda-schedule option as a replacement

ethminer/MinerAux.h

@@ -200,8 +200,19 @@ public:
 				}
 			}
 		}
-		else if (arg == "--cuda-turbo") 
-			m_cudaHighCPULoad = true;
+		else if (arg == "--cuda-schedule" && i + 1 < argc)
+		{
+			string mode = argv[++i];
+			if (mode == "Auto") m_cudaSchedule = 0;
+			else if (mode == "Spin") m_cudaSchedule = 1;
+			else if (mode == "Yield") m_cudaSchedule = 2;
+			else if (mode == "Sync") m_cudaSchedule = 4;
+			else 
+			{
+				cerr << "Bad " << arg << " option: " << argv[i] << endl;
+				BOOST_THROW_EXCEPTION(BadArgument());
+			}
+		}
 		else if (arg == "--cuda-streams" && i + 1 < argc)
 			m_numStreams = stol(argv[++i]);
 #endif
@@ -393,7 +404,7 @@ public:
 				m_globalWorkSizeMultiplier,
 				m_numStreams,
 				m_extraGPUMemory,
-				m_cudaHighCPULoad,
+				m_cudaSchedule,
 				m_currentBlock
 				))
 				exit(1);
@@ -457,7 +468,11 @@ public:
 			<< "    --cuda-block-size Set the CUDA block work size. Default is " << toString(ethash_cuda_miner::c_defaultBlockSize) << endl
 			<< "    --cuda-grid-size Set the CUDA grid size. Default is " << toString(ethash_cuda_miner::c_defaultGridSize) << endl
 			<< "    --cuda-streams Set the number of CUDA streams. Default is " << toString(ethash_cuda_miner::c_defaultNumStreams) << endl
-			<< "    --cuda-turbo Get a bit of extra hashrate at the cost of high CPU load... Default is false" << endl
+			<< "    --cuda-schedule <mode> Set the schedule mode for CUDA threads waiting for CUDA devices to finish work. Default is Auto. Possible values are:" << endl
+			<< "        Auto  - Uses a heuristic based on the number of active CUDA contexts in the process C and the number of logical processors in the system P. If C > P, then Yield else Spin." << endl
+			<< "        Spin  - Instruct CUDA to actively spin when waiting for results from the device." << endl
+			<< "        Yield - Instruct CUDA to yield its thread when waiting for results from the device." << endl
+			<< "        Sync  - Instruct CUDA to block the CPU thread on a synchronization primitive when waiting for the results from the device." << endl
 			<< "    --cuda-devices <0 1 ..n> Select which GPUs to mine on. Default is to use all" << endl
 #endif
 			;
@@ -795,7 +810,7 @@ private:
 	unsigned m_cudaDeviceCount = 0;
 	unsigned m_cudaDevices[16];
 	unsigned m_numStreams = ethash_cuda_miner::c_defaultNumStreams;
-	bool m_cudaHighCPULoad = false;
+	unsigned m_cudaSchedule = 0;
 #endif
 	uint64_t m_currentBlock = 0;
 	// default value is 350MB of GPU memory for other stuff (windows system rendering, e.t.c.)

libethash-cuda/ethash_cuda_miner.cpp

@@ -130,7 +130,7 @@ bool ethash_cuda_miner::configureGPU(
 	unsigned _gridSize,
 	unsigned _numStreams,
 	unsigned _extraGPUMemory,
-	bool	 _highcpu,
+	unsigned _scheduleFlag,
 	uint64_t _currentBlock
 	)
 {
@@ -138,7 +138,7 @@ bool ethash_cuda_miner::configureGPU(
 	s_gridSize = _gridSize;
 	s_extraRequiredGPUMem = _extraGPUMemory;
 	s_numStreams = _numStreams;
-	s_highCPU = _highcpu;
+	s_scheduleFlag = _scheduleFlag;
 
 	// by default let's only consider the DAG of the first epoch
 	uint64_t dagSize = ethash_get_datasize(_currentBlock);
@@ -174,7 +174,7 @@ unsigned ethash_cuda_miner::s_extraRequiredGPUMem;
 unsigned ethash_cuda_miner::s_blockSize = ethash_cuda_miner::c_defaultBlockSize;
 unsigned ethash_cuda_miner::s_gridSize = ethash_cuda_miner::c_defaultGridSize;
 unsigned ethash_cuda_miner::s_numStreams = ethash_cuda_miner::c_defaultNumStreams;
-bool ethash_cuda_miner::s_highCPU = false;
+unsigned ethash_cuda_miner::s_scheduleFlag = 0;
 
 void ethash_cuda_miner::listDevices()
 {
@@ -226,7 +226,7 @@ bool ethash_cuda_miner::init(uint8_t const* _dag, uint64_t _dagSize, unsigned _d
 		return false;
 	}
 	cudaDeviceReset();
-	cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
+	cudaSetDeviceFlags(s_scheduleFlag);
 	cudaDeviceSetCacheConfig(cudaFuncCachePreferL1);
 
 	m_search_buf = new uint32_t *[s_numStreams];
@@ -263,19 +263,6 @@ bool ethash_cuda_miner::init(uint8_t const* _dag, uint64_t _dagSize, unsigned _d
 	return true;
 }
 
-/**
- * Prevent High CPU usage while waiting for an async task
- */
-static unsigned waitStream(cudaStream_t stream)
-{
-	unsigned wait_ms = 0;
-	while (cudaStreamQuery(stream) == cudaErrorNotReady) {
-		this_thread::sleep_for(chrono::milliseconds(10));
-		wait_ms += 10;
-	}
-	return wait_ms;
-}
-
 void ethash_cuda_miner::search(uint8_t const* header, uint64_t target, search_hook& hook)
 {
 	struct pending_batch
@@ -313,10 +300,7 @@ void ethash_cuda_miner::search(uint8_t const* header, uint64_t target, search_ho
 		{
 			pending_batch const& batch = pending.front();
 
-			if (s_highCPU)
-				cudaStreamSynchronize(m_streams[buf]);
-			else
-				waitStream(m_streams[buf]); // 28ms
+			cudaStreamSynchronize(m_streams[buf]);
 
 			uint32_t * results = m_search_buf[batch.buf];
 			unsigned num_found = std::min<unsigned>(results[0], c_max_search_results);

libethash-cuda/ethash_cuda_miner.h

@@ -31,7 +31,7 @@ public:
 		unsigned _gridSize,
 		unsigned _numStreams,
 		unsigned _extraGPUMemory,
-		bool	 _highcpu,
+		unsigned _scheduleFlag,
 		uint64_t _currentBlock
 		);
 	bool init(
@@ -66,8 +66,8 @@ private:
 	static unsigned s_gridSize;
 	/// The number of CUDA streams
 	static unsigned s_numStreams;
-	/// Whether or not to let the CPU wait
-	static bool s_highCPU;
+	/// CUDA schedule flag
+	static unsigned s_scheduleFlag;
 
 	/// GPU memory required for other things, like window rendering e.t.c.
 	/// User can set it via the --cl-extragpu-mem argument.

libethcore/EthashCUDAMiner.cpp

@@ -215,7 +215,7 @@ bool EthashCUDAMiner::configureGPU(
 	unsigned _gridSize,
 	unsigned _numStreams,
 	unsigned _extraGPUMemory,
-	bool _highcpu,
+	unsigned _scheduleFlag,
 	uint64_t _currentBlock
 	)
 {
@@ -230,7 +230,7 @@ bool EthashCUDAMiner::configureGPU(
 		_gridSize,
 		_numStreams,
 		_extraGPUMemory,
-		_highcpu,
+		_scheduleFlag,
 		_currentBlock)
 		)
 	{

libethcore/EthashCUDAMiner.h

@@ -52,7 +52,7 @@ namespace eth
 			unsigned _gridSize,
 			unsigned _numStreams,
 			unsigned _extraGPUMemory,
-			bool	 _highcpu,
+			unsigned _scheduleFlag,
 			uint64_t _currentBlock
 			);
 		static void setNumInstances(unsigned _instances)