@ -24,6 +24,7 @@
# include <cstdio>
# include <cstdlib>
# include <fstream>
# include <iostream>
# include <assert.h>
# include <queue>
@ -35,6 +36,7 @@
# include "ethash_cl_miner_kernel.h"
# define ETHASH_BYTES 32
# define ETHASH_CL_MINIMUM_MEMORY 2000000000
// workaround lame platforms
# if !CL_VERSION_1_2
@ -47,6 +49,9 @@
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 )
{
char buf [ 256 ] ;
@ -72,7 +77,7 @@ std::string ethash_cl_miner::platform_info(unsigned _platformId, unsigned _devic
cl : : Platform : : get ( & platforms ) ;
if ( platforms . empty ( ) )
{
cout < < " No OpenCL platforms found. " < < endl ;
ETHCL_LOG ( " No OpenCL platforms found. " ) ;
return std : : string ( ) ;
}
@ -82,7 +87,7 @@ std::string ethash_cl_miner::platform_info(unsigned _platformId, unsigned _devic
platforms [ platform_num ] . getDevices ( CL_DEVICE_TYPE_ALL , & devices ) ;
if ( devices . empty ( ) )
{
cout < < " No OpenCL devices found. " < < endl ;
ETHCL_LOG ( " No OpenCL devices found. " ) ;
return std : : string ( ) ;
}
@ -107,7 +112,7 @@ unsigned ethash_cl_miner::get_num_devices(unsigned _platformId)
cl : : Platform : : get ( & platforms ) ;
if ( platforms . empty ( ) )
{
cout < < " No OpenCL platforms found. " < < endl ;
ETHCL_LOG ( " No OpenCL platforms found. " ) ;
return 0 ;
}
@ -116,52 +121,105 @@ unsigned ethash_cl_miner::get_num_devices(unsigned _platformId)
platforms [ platform_num ] . getDevices ( CL_DEVICE_TYPE_ALL , & devices ) ;
if ( devices . empty ( ) )
{
cout < < " No OpenCL devices found. " < < endl ;
ETHCL_LOG ( " No OpenCL devices found. " ) ;
return 0 ;
}
return devices . size ( ) ;
}
bool ethash_cl_miner : : haveSufficientGPUMemory ( unsigned _platformId )
{
std : : vector < cl : : Platform > platforms ;
cl : : Platform : : get ( & platforms ) ;
if ( platforms . empty ( ) )
{
ETHCL_LOG ( " No OpenCL platforms found. " ) ;
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 ( ) )
{
ETHCL_LOG ( " No OpenCL devices found. " ) ;
return false ;
}
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 "
) ;
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 : : finish ( )
{
if ( m_queue ( ) )
m_queue . finish ( ) ;
}
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_dagChunksNum = _dagChunksNum ;
// get all platforms
try
{
std : : vector < cl : : Platform > platforms ;
cl : : Platform : : get ( & platforms ) ;
if ( platforms . empty ( ) )
{
cout < < " No OpenCL platforms found. " < < endl ;
ETHCL_LOG ( " No OpenCL platforms found. " ) ;
return false ;
}
// use selected platform
_platformId = std : : min < unsigned > ( _platformId , platforms . size ( ) - 1 ) ;
cout < < " Using platform: " < < platforms [ _platformId ] . getInfo < CL_PLATFORM_NAME > ( ) . c_str ( ) < < endl ;
ETHCL_LOG ( " Using platform: " < < platforms [ _platformId ] . getInfo < CL_PLATFORM_NAME > ( ) . c_str ( ) ) ;
// get GPU device of the default platform
std : : vector < cl : : Device > devices ;
platforms [ _platformId ] . getDevices ( CL_DEVICE_TYPE_ALL , & devices ) ;
if ( devices . empty ( ) )
{
cout < < " No OpenCL devices found. " < < endl ;
ETHCL_LOG ( " No OpenCL devices found. " ) ;
return false ;
}
// use selected device
cl : : Device & device = devices [ std : : min < unsigned > ( _deviceId , devices . size ( ) - 1 ) ] ;
std : : string device_version = device . getInfo < CL_DEVICE_VERSION > ( ) ;
cout < < " Using device: " < < device . getInfo < CL_DEVICE_NAME > ( ) . c_str ( ) < < " ( " < < device_version . c_str ( ) < < " ) " < < endl ;
ETHCL_LOG ( " Using device: " < < device . getInfo < CL_DEVICE_NAME > ( ) . c_str ( ) < < " ( " < < device_version . c_str ( ) < < " ) " ) ;
if ( strncmp ( " OpenCL 1.0 " , device_version . c_str ( ) , 10 ) = = 0 )
{
cout < < " OpenCL 1.0 is not supported. " < < endl ;
ETHCL_LOG ( " OpenCL 1.0 is not supported. " ) ;
return false ;
}
if ( strncmp ( " OpenCL 1.1 " , device_version . c_str ( ) , 10 ) = = 0 )
@ -175,6 +233,8 @@ bool ethash_cl_miner::init(uint8_t const* _dag, uint64_t _dagSize, unsigned work
m_workgroup_size = ( ( workgroup_size + 7 ) / 8 ) * 8 ;
// 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 ) ) ;
@ -184,119 +244,96 @@ bool ethash_cl_miner::init(uint8_t const* _dag, uint64_t _dagSize, unsigned work
// create miner OpenCL program
cl : : Program : : Sources sources ;
sources . push_back ( { code . c_str ( ) , code . size ( ) } ) ;
sources . push_back ( { code . c_str ( ) , code . size ( ) } ) ;
cl : : Program program ( m_context , sources ) ;
try
{
program . build ( { device } ) ;
program . build ( { device } ) ;
ETHCL_LOG ( " Printing program log " ) ;
ETHCL_LOG ( program . getBuildInfo < CL_PROGRAM_BUILD_LOG > ( device ) . c_str ( ) ) ;
}
catch ( cl : : Error err )
{
cout < < program . getBuildInfo < CL_PROGRAM_BUILD_LOG > ( device ) . c_str ( ) ;
ETHCL_LOG ( program . getBuildInfo < CL_PROGRAM_BUILD_LOG > ( device ) . c_str ( ) ) ;
return false ;
}
if ( _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
m_dag = cl : : Buffer ( m_context , CL_MEM_READ_ONLY , _dagSize ) ;
if ( _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 + + )
{
// TODO Note: If we ever change to _dagChunksNum other than 4, then the size would need recalculation
ETHCL_LOG ( " Creating buffer for chunk " < < i ) ;
m_dagChunks . push_back ( cl : : Buffer (
m_context ,
CL_MEM_READ_ONLY ,
( i = = 3 ) ? ( _dagSize - 3 * ( ( _dagSize > > 9 ) < < 7 ) ) : ( _dagSize > > 9 ) < < 7
) ) ;
}
// create buffer for header
ETHCL_LOG ( " Creating buffer for header. " ) ;
m_header = cl : : Buffer ( m_context , CL_MEM_READ_ONLY , 32 ) ;
// compute dag on CPU
try {
m_queue . enqueueWriteBuffer ( m_dag , CL_TRUE , 0 , _dagSize , _dag ) ;
if ( _dagChunksNum = = 1 )
{
ETHCL_LOG ( " Mapping one big chunk. " ) ;
m_queue . enqueueWriteBuffer ( m_dagChunks [ 0 ] , CL_TRUE , 0 , _dagSize , _dag ) ;
}
else
{
// 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 + + )
{
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 ) ;
}
catch ( . . . )
for ( unsigned i = 0 ; i < _dagChunksNum ; i + + )
{
// didn't work. shitty driver. try allocating in CPU RAM and manually memcpying it.
void * dag_ptr = m_queue . enqueueMapBuffer ( m_dag , true , m_opencl_1_1 ? CL_MAP_WRITE : CL_MAP_WRITE_INVALIDATE_REGION , 0 , _dagSize ) ;
memcpy ( dag_ptr , _dag , _dagSize ) ;
m_queue . enqueueUnmapMemObject ( m_dag , dag_ptr ) ;
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 ] ) ;
}
}
// create mining buffers
for ( unsigned i = 0 ; i ! = c_num_buffers ; + + i )
{
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 ) ;
ETHCL_LOG ( " Creating mining buffer " < < i ) ;
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 ) ) ;
}
return true ;
}
void ethash_cl_miner : : hash ( uint8_t * ret , uint8_t const * header , uint64_t nonce , unsigned count )
{
struct pending_batch
{
unsigned base ;
unsigned count ;
unsigned buf ;
} ;
std : : queue < pending_batch > pending ;
// update header constant buffer
m_queue . enqueueWriteBuffer ( m_header , true , 0 , 32 , header ) ;
/*
__kernel void ethash_combined_hash (
__global hash32_t * g_hashes ,
__constant hash32_t const * g_header ,
__global hash128_t const * g_dag ,
ulong start_nonce ,
uint isolate
)
*/
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
unsigned buf = 0 ;
for ( unsigned i = 0 ; i < count | | ! pending . empty ( ) ; )
{
// how many this batch
if ( i < count )
{
unsigned const this_count = std : : min < unsigned > ( count - i , c_hash_batch_size ) ;
unsigned const batch_count = std : : max < unsigned > ( this_count , m_workgroup_size ) ;
// supply output hash buffer to kernel
m_hash_kernel . setArg ( 0 , m_hash_buf [ buf ] ) ;
// execute it!
m_queue . enqueueNDRangeKernel (
m_hash_kernel ,
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 ;
}
// read results
if ( i = = count | | pending . size ( ) = = c_num_buffers )
catch ( cl : : Error err )
{
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 ( ) ;
}
ETHCL_LOG ( err . what ( ) < < " ( " < < err . err ( ) < < " ) " ) ;
return false ;
}
return true ;
}
void ethash_cl_miner : : search ( uint8_t const * header , uint64_t target , search_hook & hook )
{
try
{
struct pending_batch
{
uint64_t start_nonce ;
@ -304,7 +341,7 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook
} ;
std : : queue < pending_batch > pending ;
uint32_t const c_zero = 0 ;
static uint32_t const c_zero = 0 ;
// update header constant buffer
m_queue . enqueueWriteBuffer ( m_header , false , 0 , 32 , header ) ;
@ -319,37 +356,30 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook
# endif
m_queue . finish ( ) ;
/*
__kernel void ethash_combined_search (
__global hash32_t * g_hashes , // 0
__constant hash32_t const * g_header , // 1
__global hash128_t const * g_dag , // 2
ulong start_nonce , // 3
ulong target , // 4
uint isolate // 5
)
*/
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 ) ;
m_search_kernel . setArg ( argPos + 1 , target ) ;
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 ) ;
for ( ; ; start_nonce + = c_search_batch_size )
for ( ; ; start_nonce + = c_search_batch_size )
{
// supply output buffer to kernel
m_search_kernel . setArg ( 0 , m_search_buf [ buf ] ) ;
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 ) ;
pending . push ( { start_nonce , buf } ) ;
pending . push ( { start_nonce , buf } ) ;
buf = ( buf + 1 ) % c_num_buffers ;
// read results
@ -358,17 +388,14 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook
pending_batch const & batch = pending . front ( ) ;
// 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 ) ) ;
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 ) ;
uint64_t nonces [ c_max_search_results ] ;
for ( unsigned i = 0 ; i ! = num_found ; + + i )
{
nonces [ i ] = batch . start_nonce + results [ i + 1 ] ;
}
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 )
@ -387,5 +414,9 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook
if ( ! m_opencl_1_1 )
pre_return_event . wait ( ) ;
# endif
}
catch ( cl : : Error err )
{
ETHCL_LOG ( err . what ( ) < < " ( " < < err . err ( ) < < " ) " ) ;
}
}
yann300
10 years ago