diff --git a/ethminer/MinerAux.h b/ethminer/MinerAux.h
index 47fd2e2ae..b6d87e181 100644
--- a/ethminer/MinerAux.h
+++ b/ethminer/MinerAux.h
@@ -127,6 +127,10 @@ public:
cerr << "Bad " << arg << " option: " << argv[i] << endl;
throw BadArgument();
}
+ else if (arg == "--use-chunks")
+ {
+ dagChunks = 4;
+ }
else if (arg == "--phone-home" && i + 1 < argc)
{
string m = argv[++i];
@@ -293,6 +297,7 @@ 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
;
}
@@ -480,6 +485,7 @@ private:
unsigned openclPlatform = 0;
unsigned openclDevice = 0;
unsigned miningThreads = UINT_MAX;
+ unsigned dagChunks = 1;
/// DAG initialisation param.
unsigned initDAG = 0;
diff --git a/libethash-cl/ethash_cl_miner.cpp b/libethash-cl/ethash_cl_miner.cpp
index 520a13180..5e69df1ee 100644
--- a/libethash-cl/ethash_cl_miner.cpp
+++ b/libethash-cl/ethash_cl_miner.cpp
@@ -43,9 +43,6 @@
#define CL_MEM_HOST_READ_ONLY 0
#endif
-// maybe move to CMakeLists.txt ?
-// #define ETHASH_CL_CHUNK_UPLOAD
-
#undef min
#undef max
@@ -61,7 +58,7 @@ static void add_definition(std::string& source, char const* id, unsigned value)
ethash_cl_miner::search_hook::~search_hook() {}
ethash_cl_miner::ethash_cl_miner()
-: m_opencl_1_1()
+: m_dagChunks(nullptr), m_opencl_1_1()
{
}
@@ -130,10 +127,26 @@ void ethash_cl_miner::finish()
{
if (m_queue())
m_queue.finish();
+
+ if (m_dagChunks)
+ delete [] m_dagChunks;
}
-bool ethash_cl_miner::init(uint8_t const* _dag, uint64_t _dagSize, unsigned workgroup_size, unsigned _platformId, unsigned _deviceId)
+bool ethash_cl_miner::init(
+ uint8_t const* _dag,
+ uint64_t _dagSize,
+ unsigned workgroup_size,
+ unsigned _platformId,
+ unsigned _deviceId,
+ unsigned _dagChunksNum
+)
{
+ // 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_dagChunks = new cl::Buffer[_dagChunksNum];
+ m_dagChunksNum = _dagChunksNum;
+
// get all platforms
try
{
@@ -205,50 +218,61 @@ bool ethash_cl_miner::init(uint8_t const* _dag, uint64_t _dagSize, unsigned work
cout << program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(device).c_str();
return false;
}
- #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");
- cout << "loading ethash_search" << endl;
- m_search_kernel = cl::Kernel(program, "ethash_search");
- #endif
-
- // create buffer for dag
- #ifdef ETHASH_CL_CHUNK_UPLOAD
- for (unsigned i = 0; i < 4; i++)
+ if (_dagChunksNum == 1)
+ {
+ cout << "loading ethash_hash" << endl;
+ m_hash_kernel = cl::Kernel(program, "ethash_hash");
+ cout << "loading ethash_search" << endl;
+ m_search_kernel = cl::Kernel(program, "ethash_search");
+ }
+ else
{
- 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);
+ 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
+
+ // create buffer for dag
+ if (_dagChunksNum == 1)
+ m_dagChunks[0] = cl::Buffer(m_context, CL_MEM_READ_ONLY, _dagSize);
+ else
+ for (unsigned i = 0; i < _dagChunksNum; i++)
+ {
+ // TODO Note: If we ever change to _dagChunksNum other than 4, then the size would need recalculation
+ cout << "Creating buffer for chunk " << i << endl;
+ m_dagChunks[i] = cl::Buffer(
+ m_context,
+ CL_MEM_READ_ONLY,
+ (i == 3) ? (_dagSize - 3 * ((_dagSize >> 9) << 7)) : (_dagSize >> 9) << 7
+ );
+ }
cout << "Creating one big buffer." << endl;
- m_dag = cl::Buffer(m_context, CL_MEM_READ_ONLY, _dagSize);
- #endif
// create buffer for header
cout << "Creating buffer for header." << endl;
m_header = cl::Buffer(m_context, CL_MEM_READ_ONLY, 32);
- #ifdef ETHASH_CL_CHUNK_UPLOAD
- void* dag_ptr[4];
- for (unsigned i = 0; i < 4; i++)
+ if (_dagChunksNum == 1)
{
- cout << "Mapping chunk " << i << endl;
- dag_ptr[i] = m_queue.enqueueMapBuffer(m_dags[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);
+ cout << "Mapping one big chunk." << endl;
+ m_queue.enqueueWriteBuffer(m_dagChunks[0], CL_TRUE, 0, _dagSize, _dag);
}
- for (unsigned i = 0; i < 4; i++)
+ else
{
- 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]);
+ // 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++)
+ {
+ cout << "Mapping chunk " << i << endl;
+ 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++)
+ {
+ 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]);
+ }
}
- #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)
@@ -288,22 +312,13 @@ void ethash_cl_miner::hash(uint8_t* ret, uint8_t const* header, uint64_t nonce,
uint isolate
)
*/
- #ifdef ETHASH_CL_CHUNK_UPLOAD
cout << "Setting chunk hash arguments." << endl;
+ unsigned argPos = 2;
m_hash_kernel.setArg(1, m_header);
- m_hash_kernel.setArg(2, m_dags[0]);
- m_hash_kernel.setArg(3, m_dags[1]);
- m_hash_kernel.setArg(4, m_dags[2]);
- m_hash_kernel.setArg(5, m_dags[3]);
- m_hash_kernel.setArg(6, nonce);
- m_hash_kernel.setArg(7, ~0u); // have to pass this to stop the compile unrolling the loop
- #else
- cout << "Setting hash arguments." << endl;
- m_hash_kernel.setArg(1, m_header);
- 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
+ for (unsigned i = 0 ; i < m_dagChunksNum; ++i, ++argPos)
+ m_hash_kernel.setArg(argPos, m_dagChunks[i]);
+ m_hash_kernel.setArg(argPos + 1, nonce);
+ m_hash_kernel.setArg(argPos + 2, ~0u); // have to pass this to stop the compiler unrolling the loop
unsigned buf = 0;
for (unsigned i = 0; i < count || !pending.empty(); )
@@ -381,27 +396,13 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook
uint isolate // 5
)
*/
- #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]);
- m_search_kernel.setArg(3, m_dags[1]);
- m_search_kernel.setArg(4, m_dags[2]);
- m_search_kernel.setArg(5, m_dags[3]);
-
- // pass these to stop the compiler unrolling the loops
- m_search_kernel.setArg(7, target);
- m_search_kernel.setArg(8, ~0u);
-
- #else
- cout << "Setting search arguments." << endl;
+ unsigned argPos = 2;
m_search_kernel.setArg(1, m_header);
- m_search_kernel.setArg(2, m_dag);
-
+ for (unsigned i = 0; i < m_dagChunksNum; ++i, ++argPos)
+ m_search_kernel.setArg(argPos, m_dagChunks[i]);
// pass these to stop the compiler unrolling the loops
- m_search_kernel.setArg(4, target);
- m_search_kernel.setArg(5, ~0u);
- #endif
+ m_search_kernel.setArg(argPos + 1, target);
+ m_search_kernel.setArg(argPos + 2, ~0u);
unsigned buf = 0;
std::random_device engine;
@@ -410,11 +411,10 @@ 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 ETHASH_CL_CHUNK_UPLOAD
- m_search_kernel.setArg(6, start_nonce);
- #else
- m_search_kernel.setArg(3, start_nonce);
- #endif
+ if (m_dagChunksNum == 1)
+ m_search_kernel.setArg(3, start_nonce);
+ else
+ m_search_kernel.setArg(6, start_nonce);
// execute it!
m_queue.enqueueNDRangeKernel(m_search_kernel, cl::NullRange, c_search_batch_size, m_workgroup_size);
diff --git a/libethash-cl/ethash_cl_miner.h b/libethash-cl/ethash_cl_miner.h
index 9c97f2aa4..0f83f8565 100644
--- a/libethash-cl/ethash_cl_miner.h
+++ b/libethash-cl/ethash_cl_miner.h
@@ -36,7 +36,14 @@ public:
static unsigned get_num_devices(unsigned _platformId = 0);
static std::string platform_info(unsigned _platformId = 0, unsigned _deviceId = 0);
- bool init(uint8_t const* _dag, uint64_t _dagSize, unsigned workgroup_size = 64, unsigned _platformId = 0, unsigned _deviceId = 0);
+ bool init(
+ uint8_t const* _dag,
+ uint64_t _dagSize,
+ unsigned workgroup_size = 64,
+ unsigned _platformId = 0,
+ unsigned _deviceId = 0,
+ unsigned _dagChunksNum = 1
+ );
void finish();
void hash(uint8_t* ret, uint8_t const* header, uint64_t nonce, unsigned count);
void search(uint8_t const* header, uint64_t target, search_hook& hook);
@@ -51,11 +58,11 @@ private:
cl::CommandQueue m_queue;
cl::Kernel m_hash_kernel;
cl::Kernel m_search_kernel;
- cl::Buffer m_dag;
- cl::Buffer m_dags[4];
+ unsigned m_dagChunksNum;
+ cl::Buffer* m_dagChunks;
cl::Buffer m_header;
cl::Buffer m_hash_buf[c_num_buffers];
cl::Buffer m_search_buf[c_num_buffers];
unsigned m_workgroup_size;
bool m_opencl_1_1;
-};
\ No newline at end of file
+};
diff --git a/libethcore/Ethash.cpp b/libethcore/Ethash.cpp
index f62c1f9cd..158f40981 100644
--- a/libethcore/Ethash.cpp
+++ b/libethcore/Ethash.cpp
@@ -285,6 +285,7 @@ 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),
@@ -345,7 +346,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);
+ m_miner->init(dagData.data(), dagData.size(), 32, s_platformId, device, s_dagChunks);
}
uint64_t upper64OfBoundary = (uint64_t)(u64)((u256)w.boundary >> 192);
diff --git a/libethcore/Ethash.h b/libethcore/Ethash.h
index 86540678f..868c27916 100644
--- a/libethcore/Ethash.h
+++ b/libethcore/Ethash.h
@@ -119,6 +119,7 @@ public:
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;
@@ -137,6 +138,7 @@ public:
static unsigned s_platformId;
static unsigned s_deviceId;
static unsigned s_numInstances;
+ static unsigned s_dagChunks;
};
#else
using GPUMiner = CPUMiner;