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

Merge pull request #2132 from LefterisJP/opencl_chunk_improve 

Refactoring around OpenCL DAG chunks usage
cl-refactor
Gav Wood 10 years ago
parent
27fbfd4715 c9371fdf36
commit
4b10d7f29b
5 changed files with 113 additions and 110 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. 8
      ethminer/MinerAux.h
  2. 182
      libethash-cl/ethash_cl_miner.cpp
  3. 18
      libethash-cl/ethash_cl_miner.h
  4. 9
      libethcore/Ethash.cpp
  5. 6
      libethcore/Ethash.h

8
ethminer/MinerAux.h
View File

@ -132,10 +132,6 @@ public:
ProofOfWork::GPUMiner::listDevices();
exit(0);
}
else if (arg == "--use-chunks")
{
dagChunks = 4;
}
else if (arg == "--phone-home" && i + 1 < argc)
{
string m = argv[++i];
@ -180,7 +176,7 @@ public:
m_minerType = MinerType::CPU;
else if (arg == "-G" || arg == "--opencl")
{
if (!ProofOfWork::GPUMiner::haveSufficientMemory())
if (!ProofOfWork::GPUMiner::configureGPU())
{
cout << "No GPU device with sufficient memory was found. Defaulting to CPU" << endl;
m_minerType = MinerType::CPU;
@ -273,7 +269,6 @@ public:
ProofOfWork::GPUMiner::setDefaultPlatform(openclPlatform);
ProofOfWork::GPUMiner::setDefaultDevice(openclDevice);
ProofOfWork::GPUMiner::setNumInstances(miningThreads);
ProofOfWork::GPUMiner::setDagChunks(dagChunks);
}
if (mode == OperationMode::DAGInit)
doInitDAG(initDAG);
@ -311,7 +306,6 @@ 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
;
}

182
libethash-cl/ethash_cl_miner.cpp
View File

@ -52,11 +52,11 @@ using namespace std;
// TODO: If at any point we can use libdevcore in here then we should switch to using a LogChannel
#define ETHCL_LOG(_contents) cout << "[OPENCL]:" << _contents << endl
static void add_definition(std::string& source, char const* id, unsigned value)
static void addDefinition(string& _source, char const* _id, unsigned _value)
{
char buf[256];
sprintf(buf, "#define %s %uu\n", id, value);
source.insert(source.begin(), buf, buf + strlen(buf));
sprintf(buf, "#define %s %uu\n", _id, _value);
_source.insert(_source.begin(), buf, buf + strlen(buf));
}
ethash_cl_miner::search_hook::~search_hook() {}
@ -71,44 +71,44 @@ ethash_cl_miner::~ethash_cl_miner()
finish();
}
std::string ethash_cl_miner::platform_info(unsigned _platformId, unsigned _deviceId)
string ethash_cl_miner::platform_info(unsigned _platformId, unsigned _deviceId)
{
std::vector<cl::Platform> platforms;
vector<cl::Platform> platforms;
cl::Platform::get(&platforms);
if (platforms.empty())
{
ETHCL_LOG("No OpenCL platforms found.");
return std::string();
return string();
}
// get GPU device of the selected platform
std::vector<cl::Device> devices;
unsigned platform_num = std::min<unsigned>(_platformId, platforms.size() - 1);
vector<cl::Device> devices;
unsigned platform_num = min<unsigned>(_platformId, platforms.size() - 1);
platforms[platform_num].getDevices(CL_DEVICE_TYPE_ALL, &devices);
if (devices.empty())
{
ETHCL_LOG("No OpenCL devices found.");
return std::string();
return string();
}
// use selected default device
unsigned device_num = std::min<unsigned>(_deviceId, devices.size() - 1);
unsigned device_num = min<unsigned>(_deviceId, devices.size() - 1);
cl::Device& device = devices[device_num];
std::string device_version = device.getInfo<CL_DEVICE_VERSION>();
string device_version = device.getInfo<CL_DEVICE_VERSION>();
return "{ \"platform\": \"" + platforms[platform_num].getInfo<CL_PLATFORM_NAME>() + "\", \"device\": \"" + device.getInfo<CL_DEVICE_NAME>() + "\", \"version\": \"" + device_version + "\" }";
}
unsigned ethash_cl_miner::get_num_platforms()
unsigned ethash_cl_miner::getNumPlatforms()
{
std::vector<cl::Platform> platforms;
vector<cl::Platform> platforms;
cl::Platform::get(&platforms);
return platforms.size();
}
unsigned ethash_cl_miner::get_num_devices(unsigned _platformId)
unsigned ethash_cl_miner::getNumDevices(unsigned _platformId)
{
std::vector<cl::Platform> platforms;
vector<cl::Platform> platforms;
cl::Platform::get(&platforms);
if (platforms.empty())
{
@ -116,8 +116,8 @@ unsigned ethash_cl_miner::get_num_devices(unsigned _platformId)
return 0;
}
std::vector<cl::Device> devices;
unsigned platform_num = std::min<unsigned>(_platformId, platforms.size() - 1);
vector<cl::Device> devices;
unsigned platform_num = min<unsigned>(_platformId, platforms.size() - 1);
platforms[platform_num].getDevices(CL_DEVICE_TYPE_ALL, &devices);
if (devices.empty())
{
@ -127,9 +127,34 @@ unsigned ethash_cl_miner::get_num_devices(unsigned _platformId)
return devices.size();
}
bool ethash_cl_miner::haveSufficientGPUMemory()
bool ethash_cl_miner::configureGPU()
{
std::vector<cl::Platform> platforms;
return searchForAllDevices([](cl::Device const _device) -> bool
{
cl_ulong result;
_device.getInfo(CL_DEVICE_GLOBAL_MEM_SIZE, &result);
if (result >= ETHASH_CL_MINIMUM_MEMORY)
{
ETHCL_LOG(
"Found suitable OpenCL device [" << _device.getInfo<CL_DEVICE_NAME>()
<< "] with " << result << " bytes of GPU memory"
);
return true;
}
ETHCL_LOG(
"OpenCL device " << _device.getInfo<CL_DEVICE_NAME>()
<< " has insufficient GPU memory." << result <<
" bytes of memory found < " << ETHASH_CL_MINIMUM_MEMORY << " bytes of memory required"
);
return false;
}
);
}
bool ethash_cl_miner::searchForAllDevices(function<bool(cl::Device const&)> _callback)
{
vector<cl::Platform> platforms;
cl::Platform::get(&platforms);
if (platforms.empty())
{
@ -137,50 +162,31 @@ bool ethash_cl_miner::haveSufficientGPUMemory()
return false;
}
for (unsigned i = 0; i < platforms.size(); ++i)
if (haveSufficientGPUMemory(i))
if (searchForAllDevices(i, _callback))
return true;
return false;
}
bool ethash_cl_miner::haveSufficientGPUMemory(unsigned _platformId)
bool ethash_cl_miner::searchForAllDevices(unsigned _platformId, function<bool(cl::Device const&)> _callback)
{
std::vector<cl::Platform> platforms;
vector<cl::Platform> platforms;
cl::Platform::get(&platforms);
if (_platformId >= platforms.size())
return false;
std::vector<cl::Device> devices;
unsigned platform_num = std::min<unsigned>(_platformId, platforms.size() - 1);
platforms[platform_num].getDevices(CL_DEVICE_TYPE_ALL, &devices);
if (devices.empty())
return false;
vector<cl::Device> devices;
platforms[_platformId].getDevices(CL_DEVICE_TYPE_ALL, &devices);
for (cl::Device const& device: devices)
{
cl_ulong result;
device.getInfo(CL_DEVICE_GLOBAL_MEM_SIZE, &result);
if (result >= ETHASH_CL_MINIMUM_MEMORY)
{
ETHCL_LOG(
"Found suitable OpenCL device [" << device.getInfo<CL_DEVICE_NAME>()
<< "] with " << result << " bytes of GPU memory"
);
if (_callback(device))
return true;
}
else
ETHCL_LOG(
"OpenCL device " << device.getInfo<CL_DEVICE_NAME>()
<< " has insufficient GPU memory." << result <<
" bytes of memory found < " << ETHASH_CL_MINIMUM_MEMORY << " bytes of memory required"
);
}
return false;
}
void ethash_cl_miner::listDevices()
void ethash_cl_miner::doForAllDevices(function<void(cl::Device const&)> _callback)
{
std::vector<cl::Platform> platforms;
vector<cl::Platform> platforms;
cl::Platform::get(&platforms);
if (platforms.empty())
{
@ -188,26 +194,32 @@ void ethash_cl_miner::listDevices()
return;
}
for (unsigned i = 0; i < platforms.size(); ++i)
listDevices(i);
doForAllDevices(i, _callback);
}
void ethash_cl_miner::listDevices(unsigned _platformId)
void ethash_cl_miner::doForAllDevices(unsigned _platformId, function<void(cl::Device const&)> _callback)
{
std::vector<cl::Platform> platforms;
vector<cl::Platform> platforms;
cl::Platform::get(&platforms);
if (_platformId >= platforms.size())
return;
std::string outString ="Listing OpenCL devices for platform " + to_string(_platformId) + "\n[deviceID] deviceName\n";
std::vector<cl::Device> devices;
vector<cl::Device> devices;
platforms[_platformId].getDevices(CL_DEVICE_TYPE_ALL, &devices);
unsigned i = 0;
std::string deviceString;
for (cl::Device const& device: devices)
{
outString += "[" + to_string(i) + "] " + device.getInfo<CL_DEVICE_NAME>() + "\n";
++i;
}
_callback(device);
}
void ethash_cl_miner::listDevices()
{
string outString ="\nListing OpenCL devices.\nFORMAT: [deviceID] deviceName\n";
unsigned int i = 0;
doForAllDevices([&outString, &i](cl::Device const _device)
{
outString += "[" + to_string(i) + "] " + _device.getInfo<CL_DEVICE_NAME>() + "\n";
++i;
}
);
ETHCL_LOG(outString);
}
@ -222,19 +234,13 @@ bool ethash_cl_miner::init(
uint64_t _dagSize,
unsigned workgroup_size,
unsigned _platformId,
unsigned _deviceId,
unsigned _dagChunksNum
unsigned _deviceId
)
{
// 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_dagChunksNum = _dagChunksNum;
// get all platforms
try
{
std::vector<cl::Platform> platforms;
vector<cl::Platform> platforms;
cl::Platform::get(&platforms);
if (platforms.empty())
{
@ -243,11 +249,11 @@ bool ethash_cl_miner::init(
}
// use selected platform
_platformId = std::min<unsigned>(_platformId, platforms.size() - 1);
_platformId = min<unsigned>(_platformId, platforms.size() - 1);
ETHCL_LOG("Using platform: " << platforms[_platformId].getInfo<CL_PLATFORM_NAME>().c_str());
// get GPU device of the default platform
std::vector<cl::Device> devices;
vector<cl::Device> devices;
platforms[_platformId].getDevices(CL_DEVICE_TYPE_ALL, &devices);
if (devices.empty())
{
@ -256,10 +262,14 @@ bool ethash_cl_miner::init(
}
// use selected device
cl::Device& device = devices[std::min<unsigned>(_deviceId, devices.size() - 1)];
std::string device_version = device.getInfo<CL_DEVICE_VERSION>();
cl::Device& device = devices[min<unsigned>(_deviceId, devices.size() - 1)];
string device_version = device.getInfo<CL_DEVICE_VERSION>();
ETHCL_LOG("Using device: " << device.getInfo<CL_DEVICE_NAME>().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);
m_dagChunksNum = result >= ETHASH_CL_MINIMUM_MEMORY ? 4 : 1;
if (strncmp("OpenCL 1.0", device_version.c_str(), 10) == 0)
{
ETHCL_LOG("OpenCL 1.0 is not supported.");
@ -269,7 +279,7 @@ bool ethash_cl_miner::init(
m_opencl_1_1 = true;
// create context
m_context = cl::Context(std::vector<cl::Device>(&device, &device + 1));
m_context = cl::Context(vector<cl::Device>(&device, &device + 1));
m_queue = cl::CommandQueue(m_context, device);
// use requested workgroup size, but we require multiple of 8
@ -278,11 +288,11 @@ bool ethash_cl_miner::init(
// patch source code
// note: ETHASH_CL_MINER_KERNEL is simply ethash_cl_miner_kernel.cl compiled
// into a byte array by bin2h.cmake. There is no need to load the file by hand in runtime
std::string code(ETHASH_CL_MINER_KERNEL, ETHASH_CL_MINER_KERNEL + ETHASH_CL_MINER_KERNEL_SIZE);
add_definition(code, "GROUP_SIZE", m_workgroup_size);
add_definition(code, "DAG_SIZE", (unsigned)(_dagSize / ETHASH_MIX_BYTES));
add_definition(code, "ACCESSES", ETHASH_ACCESSES);
add_definition(code, "MAX_OUTPUTS", c_max_search_results);
string code(ETHASH_CL_MINER_KERNEL, ETHASH_CL_MINER_KERNEL + ETHASH_CL_MINER_KERNEL_SIZE);
addDefinition(code, "GROUP_SIZE", m_workgroup_size);
addDefinition(code, "DAG_SIZE", (unsigned)(_dagSize / ETHASH_MIX_BYTES));
addDefinition(code, "ACCESSES", ETHASH_ACCESSES);
addDefinition(code, "MAX_OUTPUTS", c_max_search_results);
//debugf("%s", code.c_str());
// create miner OpenCL program
@ -301,7 +311,7 @@ bool ethash_cl_miner::init(
ETHCL_LOG(program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(device).c_str());
return false;
}
if (_dagChunksNum == 1)
if (m_dagChunksNum == 1)
{
ETHCL_LOG("Loading single big chunk kernels");
m_hash_kernel = cl::Kernel(program, "ethash_hash");
@ -315,13 +325,13 @@ bool ethash_cl_miner::init(
}
// create buffer for dag
if (_dagChunksNum == 1)
if (m_dagChunksNum == 1)
{
ETHCL_LOG("Creating one big buffer");
m_dagChunks.push_back(cl::Buffer(m_context, CL_MEM_READ_ONLY, _dagSize));
}
else
for (unsigned i = 0; i < _dagChunksNum; i++)
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
ETHCL_LOG("Creating buffer for chunk " << i);
@ -336,7 +346,7 @@ bool ethash_cl_miner::init(
ETHCL_LOG("Creating buffer for header.");
m_header = cl::Buffer(m_context, CL_MEM_READ_ONLY, 32);
if (_dagChunksNum == 1)
if (m_dagChunksNum == 1)
{
ETHCL_LOG("Mapping one big chunk.");
m_queue.enqueueWriteBuffer(m_dagChunks[0], CL_TRUE, 0, _dagSize, _dag);
@ -345,12 +355,12 @@ bool ethash_cl_miner::init(
{
// 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++)
for (unsigned i = 0; i < m_dagChunksNum; i++)
{
ETHCL_LOG("Mapping chunk " << i);
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++)
for (unsigned i = 0; i < m_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]);
@ -382,7 +392,7 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook
uint64_t start_nonce;
unsigned buf;
};
std::queue<pending_batch> pending;
queue<pending_batch> pending;
static uint32_t const c_zero = 0;
@ -408,8 +418,8 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook
m_search_kernel.setArg(argPos + 2, ~0u);
unsigned buf = 0;
std::random_device engine;
uint64_t start_nonce = std::uniform_int_distribution<uint64_t>()(engine);
random_device engine;
uint64_t start_nonce = uniform_int_distribution<uint64_t>()(engine);
for (;; start_nonce += c_search_batch_size)
{
// supply output buffer to kernel
@ -432,7 +442,7 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook
// could use pinned host pointer instead
uint32_t* results = (uint32_t*)m_queue.enqueueMapBuffer(m_search_buf[batch.buf], true, CL_MAP_READ, 0, (1 + c_max_search_results) * sizeof(uint32_t));
unsigned num_found = std::min<unsigned>(results[0], c_max_search_results);
unsigned num_found = min<unsigned>(results[0], c_max_search_results);
uint64_t nonces[c_max_search_results];
for (unsigned i = 0; i != num_found; ++i)

18
libethash-cl/ethash_cl_miner.h
View File

@ -32,21 +32,22 @@ public:
ethash_cl_miner();
~ethash_cl_miner();
static unsigned get_num_platforms();
static unsigned get_num_devices(unsigned _platformId = 0);
static bool searchForAllDevices(unsigned _platformId, std::function<bool(cl::Device const&)> _callback);
static bool searchForAllDevices(std::function<bool(cl::Device const&)> _callback);
static void doForAllDevices(unsigned _platformId, std::function<void(cl::Device const&)> _callback);
static void doForAllDevices(std::function<void(cl::Device const&)> _callback);
static unsigned getNumPlatforms();
static unsigned getNumDevices(unsigned _platformId = 0);
static std::string platform_info(unsigned _platformId = 0, unsigned _deviceId = 0);
static bool haveSufficientGPUMemory();
static bool haveSufficientGPUMemory(unsigned _platformId);
static void listDevices();
static void listDevices(unsigned _platformId);
static bool configureGPU();
bool init(
uint8_t const* _dag,
uint64_t _dagSize,
unsigned workgroup_size = 64,
unsigned _platformId = 0,
unsigned _deviceId = 0,
unsigned _dagChunksNum = 1
unsigned _deviceId = 0
);
void finish();
void search(uint8_t const* header, uint64_t target, search_hook& hook);
@ -61,11 +62,12 @@ private:
cl::CommandQueue m_queue;
cl::Kernel m_hash_kernel;
cl::Kernel m_search_kernel;
unsigned m_dagChunksNum;
unsigned int m_dagChunksNum;
std::vector<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;
};

9
libethcore/Ethash.cpp
View File

@ -285,7 +285,6 @@ 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),
@ -347,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, s_dagChunks);
m_miner->init(dagData.data(), dagData.size(), 32, s_platformId, device);
}
uint64_t upper64OfBoundary = (uint64_t)(u64)((u256)w.boundary >> 192);
@ -374,7 +373,7 @@ std::string Ethash::GPUMiner::platformInfo()
unsigned Ethash::GPUMiner::getNumDevices()
{
return ethash_cl_miner::get_num_devices(s_platformId);
return ethash_cl_miner::getNumDevices(s_platformId);
}
void Ethash::GPUMiner::listDevices()
@ -382,9 +381,9 @@ void Ethash::GPUMiner::listDevices()
return ethash_cl_miner::listDevices();
}
bool Ethash::GPUMiner::haveSufficientMemory()
bool Ethash::GPUMiner::configureGPU()
{
return ethash_cl_miner::haveSufficientGPUMemory();
return ethash_cl_miner::configureGPU();
}
#endif

6
libethcore/Ethash.h
View File

@ -91,7 +91,7 @@ public:
static void setDagChunks(unsigned) {}
static void setDefaultDevice(unsigned) {}
static void listDevices() {}
static bool haveSufficientMemory() { return false; }
static bool configureGPU() { return false; }
static void setNumInstances(unsigned _instances) { s_numInstances = std::min<unsigned>(_instances, std::thread::hardware_concurrency()); }
protected:
void kickOff() override
@ -120,11 +120,10 @@ public:
static std::string platformInfo();
static unsigned getNumDevices();
static void listDevices();
static bool haveSufficientMemory();
static bool configureGPU();
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;
@ -143,7 +142,6 @@ public:
static unsigned s_platformId;
static unsigned s_deviceId;
static unsigned s_numInstances;
static unsigned s_dagChunks;
};
#else
using GPUMiner = CPUMiner;

Write Preview
Loading…
Cancel
Save