From 44af17201c5e2e9533dd752bfb8eff8b1d15dff6 Mon Sep 17 00:00:00 2001 From: SergiySW Date: Fri, 11 Aug 2017 18:20:15 +0300 Subject: [PATCH 1/3] long long isn't part of OpenCL --- rai/node/openclwork.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/rai/node/openclwork.cpp b/rai/node/openclwork.cpp index 21acf3ee..43ab5d24 100644 --- a/rai/node/openclwork.cpp +++ b/rai/node/openclwork.cpp @@ -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] = @@ -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]); From dc9b4c403398d320b39c9e81765f5895ad815249 Mon Sep 17 00:00:00 2001 From: SergiySW Date: Thu, 5 Oct 2017 23:11:24 +0300 Subject: [PATCH 2/3] Fixing Nvidia OpenCL --- rai/node/openclwork.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/rai/node/openclwork.cpp b/rai/node/openclwork.cpp index 43ab5d24..d81b9b60 100644 --- a/rai/node/openclwork.cpp +++ b/rai/node/openclwork.cpp @@ -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]; From c6a5a809768a881b6566d26408a8ff6540df14d5 Mon Sep 17 00:00:00 2001 From: SergiySW Date: Wed, 11 Oct 2017 16:24:00 +0300 Subject: [PATCH 3/3] Adding debug_opencl command --- rai/rai_node/entry.cpp | 90 +++++++++++++++++++++++++++++++++++++++++- 1 file changed, 88 insertions(+), 2 deletions(-) diff --git a/rai/rai_node/entry.cpp b/rai/rai_node/entry.cpp index 56c3e3d5..785e2453 100644 --- a/rai/rai_node/entry.cpp +++ b/rai/rai_node/entry.cpp @@ -5,6 +5,7 @@ #include #include +#include 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 (), "Defines the for OpenCL commands") + ("device", boost::program_options::value (), "Defines for OpenCL command") + ("threads", boost::program_options::value (), "Defines 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 (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 (vm ["platform"].as ()); + } + 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 (vm ["device"].as ()); + } + 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 (vm ["threads"].as ()); + } + 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 ::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 (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 ::max (), nullptr);