diff --git a/libethash-cl/ethash_cl_miner.cpp b/libethash-cl/ethash_cl_miner.cpp index be4e7532c..520a13180 100644 --- a/libethash-cl/ethash_cl_miner.cpp +++ b/libethash-cl/ethash_cl_miner.cpp @@ -43,13 +43,12 @@ #define CL_MEM_HOST_READ_ONLY 0 #endif -//#define CHUNKS +// maybe move to CMakeLists.txt ? +// #define ETHASH_CL_CHUNK_UPLOAD #undef min #undef max -//#define CHUNKS - using namespace std; static void add_definition(std::string& source, char const* id, unsigned value) @@ -136,7 +135,8 @@ void ethash_cl_miner::finish() bool ethash_cl_miner::init(uint8_t const* _dag, uint64_t _dagSize, unsigned workgroup_size, unsigned _platformId, unsigned _deviceId) { // get all platforms - try { + try + { std::vector platforms; cl::Platform::get(&platforms); if (platforms.empty()) @@ -146,7 +146,6 @@ bool ethash_cl_miner::init(uint8_t const* _dag, uint64_t _dagSize, unsigned work } // use selected platform - _platformId = std::min(_platformId, platforms.size() - 1); cout << "Using platform: " << platforms[_platformId].getInfo().c_str() << endl; @@ -206,12 +205,11 @@ bool ethash_cl_miner::init(uint8_t const* _dag, uint64_t _dagSize, unsigned work cout << program.getBuildInfo(device).c_str(); return false; } - #ifdef CHUNKS + #ifdef ETHASH_CL_CHUNK_UPLOAD cout << "loading ethash_hash_chunks" << endl; m_hash_kernel = cl::Kernel(program, "ethash_hash_chunks"); cout << "loading ethash_search_chunks" << endl; m_search_kernel = cl::Kernel(program, "ethash_search_chunks"); - #else cout << "loading ethash_hash" << endl; m_hash_kernel = cl::Kernel(program, "ethash_hash"); @@ -220,10 +218,10 @@ bool ethash_cl_miner::init(uint8_t const* _dag, uint64_t _dagSize, unsigned work #endif // create buffer for dag - #ifdef CHUNKS - for (unsigned i = 0; i < 4; i++){ - - cout << "Creating chunky buffer: " << i << endl; + #ifdef ETHASH_CL_CHUNK_UPLOAD + for (unsigned i = 0; i < 4; i++) + { + cout << "Creating chunky buffer: " << i << endl; m_dags[i] = cl::Buffer(m_context, CL_MEM_READ_ONLY, (i == 3) ? (_dagSize - 3 * ((_dagSize >> 9) << 7)) : (_dagSize >> 9) << 7); } #else @@ -235,7 +233,7 @@ bool ethash_cl_miner::init(uint8_t const* _dag, uint64_t _dagSize, unsigned work cout << "Creating buffer for header." << endl; m_header = cl::Buffer(m_context, CL_MEM_READ_ONLY, 32); - #ifdef CHUNKS + #ifdef ETHASH_CL_CHUNK_UPLOAD void* dag_ptr[4]; for (unsigned i = 0; i < 4; i++) { @@ -247,36 +245,25 @@ bool ethash_cl_miner::init(uint8_t const* _dag, uint64_t _dagSize, unsigned work memcpy(dag_ptr[i], (char *)_dag + i*((_dagSize >> 9) << 7), (i == 3) ? (_dagSize - 3 * ((_dagSize >> 9) << 7)) : (_dagSize >> 9) << 7); m_queue.enqueueUnmapMemObject(m_dags[i], dag_ptr[i]); } - #else - cout << "Mapping chunk." << endl; - m_queue.enqueueWriteBuffer(m_dag, CL_TRUE, 0, _dagSize, _dag); - #endif - // 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, _dagSize); - // memcpying 1GB: horrible... really. horrible. but necessary since we can't mmap *and* gpumap. - // _fillDAG(dag_ptr); - // m_queue.enqueueUnmapMemObject(m_dag, dag_ptr); - }*/ + #else + cout << "Mapping chunk." << endl; + m_queue.enqueueWriteBuffer(m_dag, CL_TRUE, 0, _dagSize, _dag); + #endif // create mining buffers for (unsigned i = 0; i != c_num_buffers; ++i) { - cout << "Creating minig buffer " << i < pending; - + // update header constant buffer m_queue.enqueueWriteBuffer(m_header, true, 0, 32, header); @@ -301,8 +288,8 @@ void ethash_cl_miner::hash(uint8_t* ret, uint8_t const* header, uint64_t nonce, uint isolate ) */ - #ifdef CHUNKS - cout << "Setting chunk hash arguments." << endl; + #ifdef ETHASH_CL_CHUNK_UPLOAD + cout << "Setting chunk hash arguments." << endl; m_hash_kernel.setArg(1, m_header); m_hash_kernel.setArg(2, m_dags[0]); m_hash_kernel.setArg(3, m_dags[1]); @@ -316,7 +303,7 @@ void ethash_cl_miner::hash(uint8_t* ret, uint8_t const* header, uint64_t nonce, m_hash_kernel.setArg(2, m_dag); m_hash_kernel.setArg(3, nonce); m_hash_kernel.setArg(4, ~0u); // have to pass this to stop the compile unrolling the loop - #endif + #endif unsigned buf = 0; for (unsigned i = 0; i < count || !pending.empty(); ) @@ -336,9 +323,9 @@ void ethash_cl_miner::hash(uint8_t* ret, uint8_t const* header, uint64_t nonce, cl::NullRange, cl::NDRange(batch_count), cl::NDRange(m_workgroup_size) - ); + ); m_queue.flush(); - + pending.push({i, this_count, buf}); i += this_count; buf = (buf + 1) % c_num_buffers; @@ -348,12 +335,10 @@ void ethash_cl_miner::hash(uint8_t* ret, uint8_t const* header, uint64_t nonce, if (i == count || pending.size() == c_num_buffers) { pending_batch const& batch = pending.front(); - // could use pinned host pointer instead, but this path isn't that important. uint8_t* hashes = (uint8_t*)m_queue.enqueueMapBuffer(m_hash_buf[batch.buf], true, CL_MAP_READ, 0, batch.count * ETHASH_BYTES); memcpy(ret + batch.base*ETHASH_BYTES, hashes, batch.count*ETHASH_BYTES); m_queue.enqueueUnmapMemObject(m_hash_buf[batch.buf], hashes); - pending.pop(); } } @@ -362,7 +347,8 @@ void ethash_cl_miner::hash(uint8_t* ret, uint8_t const* header, uint64_t nonce, void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook& hook) { - try { + try + { struct pending_batch { uint64_t start_nonce; @@ -395,7 +381,7 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook uint isolate // 5 ) */ - #ifdef CHUNKS + #ifdef ETHASH_CL_CHUNK_UPLOAD cout << "Setting chunk search arguments." << endl; m_search_kernel.setArg(1, m_header); m_search_kernel.setArg(2, m_dags[0]); @@ -407,8 +393,8 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook m_search_kernel.setArg(7, target); m_search_kernel.setArg(8, ~0u); - #else - cout << "Setting search arguments." << endl; + #else + cout << "Setting search arguments." << endl; m_search_kernel.setArg(1, m_header); m_search_kernel.setArg(2, m_dag); @@ -417,8 +403,6 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook m_search_kernel.setArg(5, ~0u); #endif - - unsigned buf = 0; std::random_device engine; uint64_t start_nonce = std::uniform_int_distribution()(engine); @@ -426,7 +410,7 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook { // supply output buffer to kernel m_search_kernel.setArg(0, m_search_buf[buf]); - #ifdef CHUNKS + #ifdef ETHASH_CL_CHUNK_UPLOAD m_search_kernel.setArg(6, start_nonce); #else m_search_kernel.setArg(3, start_nonce); @@ -449,12 +433,9 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook uint64_t nonces[c_max_search_results]; for (unsigned i = 0; i != num_found; ++i) - { nonces[i] = batch.start_nonce + results[i + 1]; - } m_queue.enqueueUnmapMemObject(m_search_buf[batch.buf], results); - bool exit = num_found && hook.found(nonces, num_found); exit |= hook.searched(batch.start_nonce, c_search_batch_size); // always report searched before exit if (exit)