lukechilds
/
ethminer
mirror of https://github.com/lukechilds/ethminer.git
1
0
Fork 0
Code Issues Projects Releases Wiki Activity
Browse Source

Style changes

cl-refactor
Lefteris Karapetsas 10 years ago
parent
8ca0277e4e
commit
aeb49b8091
1 changed files with 28 additions and 47 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Split View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 75
      libethash-cl/ethash_cl_miner.cpp

75
libethash-cl/ethash_cl_miner.cpp
View File

@ -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<cl::Platform> 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<unsigned>(_platformId, platforms.size() - 1);
cout << "Using platform: " << platforms[_platformId].getInfo<CL_PLATFORM_NAME>().c_str() << endl;
@ -206,12 +205,11 @@ 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 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 <<endl;
cout << "Creating mining buffer " << i <<endl;
m_hash_buf[i] = cl::Buffer(m_context, CL_MEM_WRITE_ONLY | (!m_opencl_1_1 ? CL_MEM_HOST_READ_ONLY : 0), 32 * c_hash_batch_size);
m_search_buf[i] = cl::Buffer(m_context, CL_MEM_WRITE_ONLY, (c_max_search_results + 1) * sizeof(uint32_t));
}
}
catch (cl::Error err)
{
std::cout << err.what() << "(" << err.err() << ")" << std::endl;
}
return true;
return true;
}
void ethash_cl_miner::hash(uint8_t* ret, uint8_t const* header, uint64_t nonce, unsigned count)
@ -288,7 +275,7 @@ void ethash_cl_miner::hash(uint8_t* ret, uint8_t const* header, uint64_t nonce,
unsigned buf;
};
std::queue<pending_batch> 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<uint64_t>()(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)

Write Preview
Loading…
Cancel
Save