Minecon724/dncurrency
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages 1 Wiki Activity Actions

Merge branch 'pulls/167'

Browse source
This commit is contained in:
clemahieu 2017-10-28 18:28:20 -05:00
commit 8d06ae468e
2 changed files with 94 additions and 8 deletions
Download patch file Download diff file Expand all files Collapse all files

12
rai/node/openclwork.cpp
View file

@ -49,10 +49,10 @@ typedef struct __blake2b_state
__constant static ulong blake2b_IV[8] =
{
0x6a09e667f3bcc908ULL, 0xbb67ae8584caa73bULL,
0x3c6ef372fe94f82bULL, 0xa54ff53a5f1d36f1ULL,
0x510e527fade682d1ULL, 0x9b05688c2b3e6c1fULL,
0x1f83d9abfb41bd6bULL, 0x5be0cd19137e2179ULL
0x6a09e667f3bcc908UL, 0xbb67ae8584caa73bUL,
0x3c6ef372fe94f82bUL, 0xa54ff53a5f1d36f1UL,
0x510e527fade682d1UL, 0x9b05688c2b3e6c1fUL,
0x1f83d9abfb41bd6bUL, 0x5be0cd19137e2179UL
};
__constant static uchar blake2b_sigma[12][16] =
@ -206,7 +206,7 @@ static inline int blake2b_init( blake2b_state *S, const uchar outlen )
return blake2b_init_param( S, P );
}
static int blake2b_compress( blake2b_state *S, const uchar block[BLAKE2B_BLOCKBYTES] )
static int blake2b_compress( blake2b_state *S, __private const uchar block[BLAKE2B_BLOCKBYTES] )
{
ulong m[16];
ulong v[16];
@ -382,7 +382,7 @@ __kernel void raiblocks_work (__global ulong * attempt, __global ulong * result_
void printstate (blake2b_state * S)
{
printf ("%llu %llu %llu %llu %llu %llu %llu %llu %llu %llu %llu %llu\n", S->h[0], S->h[1], S->h[2], S->h[3], S->h[4], S->h[5], S->h[6], S->h[7], S->t[0], S->t[1], S->f[0], S->f[1]);
printf ("%lu %lu %lu %lu %lu %lu %lu %lu %lu %lu %lu %lu\n", S->h[0], S->h[1], S->h[2], S->h[3], S->h[4], S->h[5], S->h[6], S->h[7], S->t[0], S->t[1], S->f[0], S->f[1]);
for (int i = 0; i < 256; ++i)
{
printf ("%02x", S->buf[i]);

90
rai/rai_node/entry.cpp
View file

@ -5,6 +5,7 @@
#include <argon2.h>
#include <boost/program_options.hpp>
#include <boost/lexical_cast.hpp>
class xorshift128
{
@ -110,17 +111,21 @@ int main (int argc, char * const * argv)
rai::add_node_options (description);
description.add_options ()
("help", "Print out options")
("daemon", "Start node daemon")
("daemon", "Start node daemon")
("debug_block_count", "Display the number of block")
("debug_bootstrap_generate", "Generate bootstrap sequence of blocks")
("debug_dump_representatives", "List representatives and weights")
("debug_frontier_count", "Display the number of accounts")
("debug_mass_activity", "Generates fake debug activity")
("debug_profile_generate", "Profile work generation")
("debug_opencl", "OpenCL work generation")
("debug_profile_verify", "Profile work verification")
("debug_profile_kdf", "Profile kdf function")
("debug_verify_profile", "Profile signature verification")
("debug_xorshift_profile", "Profile xorshift algorithms");
("debug_xorshift_profile", "Profile xorshift algorithms")
("platform", boost::program_options::value <std::string> (), "Defines the <platform> for OpenCL commands")
("device", boost::program_options::value <std::string> (), "Defines <device> for OpenCL command")
("threads", boost::program_options::value <std::string> (), "Defines <threads> count for OpenCL command");
boost::program_options::variables_map vm;
boost::program_options::store (boost::program_options::parse_command_line(argc, argv, description), vm);
boost::program_options::notify (vm);
@ -263,6 +268,87 @@ int main (int argc, char * const * argv)
std::cerr << boost::str (boost::format ("%|1$ 12d|\n") % std::chrono::duration_cast <std::chrono::microseconds> (end1 - begin1).count ());
}
}
else if (vm.count ("debug_opencl"))
{
bool error (false);
rai::opencl_environment environment (error);
if (!error)
{
unsigned short platform (0);
if (vm.count ("platform") == 1)
{
try {
platform = boost::lexical_cast <unsigned short> (vm ["platform"].as <std::string> ());
}
catch (boost::bad_lexical_cast & e ) {
std::cerr << "Invalid platform id\n";
result = -1;
}
}
unsigned short device (0);
if (vm.count ("device") == 1)
{
try {
device = boost::lexical_cast <unsigned short> (vm ["device"].as <std::string> ());
}
catch (boost::bad_lexical_cast & e ) {
std::cerr << "Invalid device id\n";
result = -1;
}
}
unsigned threads (1024 * 1024);
if (vm.count ("threads") == 1)
{
try {
threads = boost::lexical_cast <unsigned> (vm ["threads"].as <std::string> ());
}
catch (boost::bad_lexical_cast & e ) {
std::cerr << "Invalid threads count\n";
result = -1;
}
}
if (!result)
{
error |= platform >= environment.platforms.size ();
if (!error)
{
error |= device >= environment.platforms[platform].devices.size ();
if (!error)
{
rai::logging logging;
logging.init (rai::unique_path ());
auto work (rai::opencl_work::create (true, {platform, device, threads}, logging));
rai::work_pool work_pool (std::numeric_limits <unsigned>::max (), std::move (work));
rai::change_block block (0, 0, rai::keypair ().prv, 0, 0);
std::cerr << boost::str (boost::format ("Starting OpenCL generation profiling. Platform: %1%. Device: %2%. Threads: %3%\n") % platform % device % threads);
for (uint64_t i (0); true; ++i)
{
block.hashables.previous.qwords [0] += 1;
auto begin1 (std::chrono::high_resolution_clock::now ());
block.block_work_set (work_pool.generate (block.root ()));
auto end1 (std::chrono::high_resolution_clock::now ());
std::cerr << boost::str (boost::format ("%|1$ 12d|\n") % std::chrono::duration_cast <std::chrono::microseconds> (end1 - begin1).count ());
}
}
else
{
std::cout << "Not available device id\n" << std::endl;
result = -1;
}
}
else
{
std::cout << "Not available platform id\n" << std::endl;
result = -1;
}
}
}
else
{
std::cout << "Error initializing OpenCL" << std::endl;
result = -1;
}
}
else if (vm.count ("debug_profile_verify"))
{
rai::work_pool work (std::numeric_limits <unsigned>::max (), nullptr);