From 3302539a11a646912ecbacec8db49450d30d6d1b Mon Sep 17 00:00:00 2001 From: Lefteris Karapetsas Date: Thu, 18 Jun 2015 10:05:41 +0200 Subject: [PATCH 1/3] 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 | 11 ++---- libethash-cl/ethash_cl_miner.cpp | 61 ++++++++++++++------------------ libethash-cl/ethash_cl_miner.h | 3 -- libethcore/Ethash.cpp | 3 +- libethcore/Ethash.h | 3 +- 5 files changed, 31 insertions(+), 50 deletions(-) diff --git a/ethminer/MinerAux.h b/ethminer/MinerAux.h index 3351b90de..a609754dd 100644 --- a/ethminer/MinerAux.h +++ b/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 )) { @@ -318,10 +315,9 @@ public: << " --opencl-device When mining using -G/--opencl use OpenCL device n (default: 0)." << endl << " -t, --mining-threads Limit number of CPU/GPU miners to n (default: use everything available on selected platform)" << endl << " --allow-opencl-cpu Allows CPU to be considered as an OpenCL device if the OpenCL platform supports it." << endl - << " --list-devices List the detected OpenCL devices and exit." < m_currentBlock; // default value is 350MB of GPU memory for other stuff (windows system rendering, e.t.c.) unsigned m_extraGPUMemory = 350000000; diff --git a/libethash-cl/ethash_cl_miner.cpp b/libethash-cl/ethash_cl_miner.cpp index b160cdd94..2fc6102fb 100644 --- a/libethash-cl/ethash_cl_miner.cpp +++ b/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 _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 _callback) @@ -288,23 +285,6 @@ bool ethash_cl_miner::init( string device_version = device.getInfo(); ETHCL_LOG("Using device: " << device.getInfo().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(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."); diff --git a/libethash-cl/ethash_cl_miner.h b/libethash-cl/ethash_cl_miner.h index cc01b0057..f36082a5a 100644 --- a/libethash-cl/ethash_cl_miner.h +++ b/libethash-cl/ethash_cl_miner.h @@ -44,7 +44,6 @@ public: static bool configureGPU( bool _allowCPU, unsigned _extraGPUMemory, - bool _forceSingleChunk, boost::optional _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. diff --git a/libethcore/Ethash.cpp b/libethcore/Ethash.cpp index ebf8c5615..b277e3c1c 100644 --- a/libethcore/Ethash.cpp +++ b/libethcore/Ethash.cpp @@ -389,13 +389,12 @@ bool Ethash::GPUMiner::configureGPU( unsigned _deviceId, bool _allowCPU, unsigned _extraGPUMemory, - bool _forceSingleChunk, boost::optional _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 diff --git a/libethcore/Ethash.h b/libethcore/Ethash.h index a5a7856f1..11e012df5 100644 --- a/libethcore/Ethash.h +++ b/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) { return false; } + static bool configureGPU(unsigned, unsigned, bool, unsigned, boost::optional) { return false; } static void setNumInstances(unsigned _instances) { s_numInstances = std::min(_instances, std::thread::hardware_concurrency()); } protected: void kickOff() override @@ -122,7 +122,6 @@ public: unsigned _deviceId, bool _allowCPU, unsigned _extraGPUMemory, - bool _forceSingleChunk, boost::optional _currentBlock ); static void setNumInstances(unsigned _instances) { s_numInstances = std::min(_instances, getNumDevices()); } From 25d2fa1607a1bf8508ce3458bde7995366683805 Mon Sep 17 00:00:00 2001 From: Lefteris Karapetsas Date: Thu, 18 Jun 2015 10:38:27 +0200 Subject: [PATCH 2/3] single chunk test: catch only cl::Error --- libethash-cl/ethash_cl_miner.cpp | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/libethash-cl/ethash_cl_miner.cpp b/libethash-cl/ethash_cl_miner.cpp index 2fc6102fb..c5f77045b 100644 --- a/libethash-cl/ethash_cl_miner.cpp +++ b/libethash-cl/ethash_cl_miner.cpp @@ -334,8 +334,11 @@ bool ethash_cl_miner::init( m_dagChunks.push_back(cl::Buffer(m_context, CL_MEM_READ_ONLY, _dagSize)); ETHCL_LOG("Created one big buffer for the DAG"); } - catch (...) + catch (cl::Error err) { + int errCode = err.err(); + if (errCode != CL_INVALID_BUFFER_SIZE || errCode != CL_MEM_OBJECT_ALLOCATION_FAILURE) + ETHCL_LOG("Allocating single buffer failed with: " << err.what() << "(" << errCode << ")"); cl_ulong result; device.getInfo(CL_DEVICE_MAX_MEM_ALLOC_SIZE, &result); ETHCL_LOG( From 8e5e2f4a9e177c05e11f288993e085a8213d1dd4 Mon Sep 17 00:00:00 2001 From: Lefteris Karapetsas Date: Thu, 18 Jun 2015 10:56:47 +0200 Subject: [PATCH 3/3] Catch OpenCL exceptions by const& --- libethash-cl/ethash_cl_miner.cpp | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/libethash-cl/ethash_cl_miner.cpp b/libethash-cl/ethash_cl_miner.cpp index c5f77045b..315f29685 100644 --- a/libethash-cl/ethash_cl_miner.cpp +++ b/libethash-cl/ethash_cl_miner.cpp @@ -321,7 +321,7 @@ bool ethash_cl_miner::init( ETHCL_LOG("Printing program log"); ETHCL_LOG(program.getBuildInfo(device).c_str()); } - catch (cl::Error err) + catch (cl::Error const& err) { ETHCL_LOG(program.getBuildInfo(device).c_str()); return false; @@ -334,7 +334,7 @@ bool ethash_cl_miner::init( m_dagChunks.push_back(cl::Buffer(m_context, CL_MEM_READ_ONLY, _dagSize)); ETHCL_LOG("Created one big buffer for the DAG"); } - catch (cl::Error err) + catch (cl::Error const& err) { int errCode = err.err(); if (errCode != CL_INVALID_BUFFER_SIZE || errCode != CL_MEM_OBJECT_ALLOCATION_FAILURE) @@ -345,6 +345,7 @@ bool ethash_cl_miner::init( "Failed to allocate 1 big chunk. Max allocateable memory is " << result << ". Trying to allocate 4 chunks." ); + // The OpenCL kernel has a hard coded number of 4 chunks at the moment m_dagChunksNum = 4; for (unsigned i = 0; i < m_dagChunksNum; i++) { @@ -404,7 +405,7 @@ bool ethash_cl_miner::init( m_search_buf[i] = cl::Buffer(m_context, CL_MEM_WRITE_ONLY, (c_max_search_results + 1) * sizeof(uint32_t)); } } - catch (cl::Error err) + catch (cl::Error const& err) { ETHCL_LOG(err.what() << "(" << err.err() << ")"); return false; @@ -498,7 +499,7 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook pre_return_event.wait(); #endif } - catch (cl::Error err) + catch (cl::Error const& err) { ETHCL_LOG(err.what() << "(" << err.err() << ")"); }