diff --git a/Dispatcher.cpp b/Dispatcher.cpp index c69636e..b71c336 100644 --- a/Dispatcher.cpp +++ b/Dispatcher.cpp @@ -188,9 +188,7 @@ Dispatcher::Device::Device(Dispatcher & parent, cl_context & clContext, cl_progr m_clQueue(createQueue(clContext, clDeviceId) ), m_kernelInit( createKernel(clProgram, "profanity_init") ), m_kernelInverse(createKernel(clProgram, "profanity_inverse")), - m_kernelIterate(createKernel(clProgram, "profanity_iterate")), - m_kernelTransform( mode.transformKernel() == "" ? NULL : createKernel(clProgram, mode.transformKernel())), - m_kernelScore(createKernel(clProgram, mode.kernel)), + m_kernelIterate(createKernel(clProgram, mode.kernel)), m_memPrecomp(clContext, m_clQueue, CL_MEM_READ_ONLY | CL_MEM_HOST_WRITE_ONLY, sizeof(g_precomp), g_precomp), m_memPointsDeltaX(clContext, m_clQueue, CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS, size, true), m_memInversedNegativeDoubleGy(clContext, m_clQueue, CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS, size, true), @@ -325,23 +323,16 @@ void Dispatcher::initBegin(Device & d) { d.m_memPointsDeltaX.setKernelArg(d.m_kernelInverse, 0); d.m_memInversedNegativeDoubleGy.setKernelArg(d.m_kernelInverse, 1); - // Kernel arguments - profanity_iterate + // Kernel arguments - profanity_iterate_score_* d.m_memPointsDeltaX.setKernelArg(d.m_kernelIterate, 0); d.m_memInversedNegativeDoubleGy.setKernelArg(d.m_kernelIterate, 1); d.m_memPrevLambda.setKernelArg(d.m_kernelIterate, 2); + d.m_memResult.setKernelArg(d.m_kernelIterate, 3); + d.m_memData1.setKernelArg(d.m_kernelIterate, 4); + d.m_memData2.setKernelArg(d.m_kernelIterate, 5); - // Kernel arguments - profanity_transform_* - if(d.m_kernelTransform) { - d.m_memInversedNegativeDoubleGy.setKernelArg(d.m_kernelTransform, 0); - } - - // Kernel arguments - profanity_score_* - d.m_memInversedNegativeDoubleGy.setKernelArg(d.m_kernelScore, 0); - d.m_memResult.setKernelArg(d.m_kernelScore, 1); - d.m_memData1.setKernelArg(d.m_kernelScore, 2); - d.m_memData2.setKernelArg(d.m_kernelScore, 3); - - CLMemory::setKernelArg(d.m_kernelScore, 4, d.m_clScoreMax); // Updated in handleResult() + CLMemory::setKernelArg(d.m_kernelIterate, 6, d.m_clScoreMax); // Updated in handleResult() + CLMemory::setKernelArg(d.m_kernelIterate, 7, (cl_uchar) (m_mode.target == CONTRACT ? 1 : 0)); // Seed device initContinue(d); @@ -436,11 +427,6 @@ void Dispatcher::dispatch(Device & d) { enqueueKernelDevice(d, d.m_kernelIterate, m_size); #endif - if (d.m_kernelTransform) { - enqueueKernelDevice(d, d.m_kernelTransform, m_size); - } - - enqueueKernelDevice(d, d.m_kernelScore, m_size); clFlush(d.m_clQueue); #ifdef PROFANITY_DEBUG @@ -490,7 +476,7 @@ void Dispatcher::handleResult(Device & d) { if (r.found > 0 && i >= d.m_clScoreMax) { d.m_clScoreMax = i; - CLMemory::setKernelArg(d.m_kernelScore, 4, d.m_clScoreMax); + CLMemory::setKernelArg(d.m_kernelIterate, 6, d.m_clScoreMax); std::lock_guard lock(m_mutex); if (i >= m_clScoreMax) { diff --git a/Dispatcher.hpp b/Dispatcher.hpp index 71638af..1996270 100644 --- a/Dispatcher.hpp +++ b/Dispatcher.hpp @@ -51,8 +51,6 @@ class Dispatcher { cl_kernel m_kernelInit; cl_kernel m_kernelInverse; cl_kernel m_kernelIterate; - cl_kernel m_kernelTransform; - cl_kernel m_kernelScore; CLMemory m_memPrecomp; CLMemory m_memPointsDeltaX; diff --git a/Mode.cpp b/Mode.cpp index c6e3194..e683fbe 100644 --- a/Mode.cpp +++ b/Mode.cpp @@ -8,7 +8,7 @@ Mode::Mode() : score(0) { Mode Mode::benchmark() { Mode r; r.name = "benchmark"; - r.kernel = "profanity_score_benchmark"; + r.kernel = "profanity_iterate_score_benchmark"; return r; } @@ -40,7 +40,7 @@ static std::string::size_type hexValue(char c) { Mode Mode::matching(const std::string strHex) { Mode r; r.name = "matching"; - r.kernel = "profanity_score_matching"; + r.kernel = "profanity_iterate_score_matching"; if (strHex.size() > 40) { throw std::runtime_error("hex mask must be at most 40 characters, got " + std::to_string(strHex.size())); @@ -73,7 +73,7 @@ Mode Mode::matching(const std::string strHex) { Mode Mode::exact(const std::string strHex) { Mode r = matching(strHex); r.name = "exact"; - r.kernel = "profanity_exact_match"; + r.kernel = "profanity_iterate_exact_match"; return r; } @@ -81,7 +81,7 @@ Mode Mode::leading(const char charLeading) { Mode r; r.name = "leading"; - r.kernel = "profanity_score_leading"; + r.kernel = "profanity_iterate_score_leading"; r.data1[0] = static_cast(hexValue(charLeading)); return r; } @@ -89,7 +89,7 @@ Mode Mode::leading(const char charLeading) { Mode Mode::range(const cl_uchar min, const cl_uchar max) { Mode r; r.name = "range"; - r.kernel = "profanity_score_range"; + r.kernel = "profanity_iterate_score_range"; r.data1[0] = min; r.data2[0] = max; return r; @@ -98,7 +98,7 @@ Mode Mode::range(const cl_uchar min, const cl_uchar max) { Mode Mode::zeroBytes() { Mode r; r.name = "zeroBytes"; - r.kernel = "profanity_score_zerobytes"; + r.kernel = "profanity_iterate_score_zerobytes"; return r; } @@ -114,17 +114,6 @@ Mode Mode::numbers() { return r; } -std::string Mode::transformKernel() const { - switch (this->target) { - case ADDRESS: - return ""; - case CONTRACT: - return "profanity_transform_contract"; - default: - throw "No kernel for target"; - } -} - std::string Mode::transformName() const { switch (this->target) { case ADDRESS: @@ -139,7 +128,7 @@ std::string Mode::transformName() const { Mode Mode::leadingRange(const cl_uchar min, const cl_uchar max) { Mode r; r.name = "leadingrange"; - r.kernel = "profanity_score_leadingrange"; + r.kernel = "profanity_iterate_score_leadingrange"; r.data1[0] = min; r.data2[0] = max; return r; @@ -148,13 +137,13 @@ Mode Mode::leadingRange(const cl_uchar min, const cl_uchar max) { Mode Mode::mirror() { Mode r; r.name = "mirror"; - r.kernel = "profanity_score_mirror"; + r.kernel = "profanity_iterate_score_mirror"; return r; } Mode Mode::doubles() { Mode r; r.name = "doubles"; - r.kernel = "profanity_score_doubles"; + r.kernel = "profanity_iterate_score_doubles"; return r; } diff --git a/Mode.hpp b/Mode.hpp index 508fecb..9a412cc 100644 --- a/Mode.hpp +++ b/Mode.hpp @@ -39,8 +39,6 @@ class Mode { std::string kernel; HashTarget target; - // kernel transform fn name - std::string transformKernel() const; // Address, Contract, ... std::string transformName() const; diff --git a/README.md b/README.md index 9f36135..70396d0 100644 --- a/README.md +++ b/README.md @@ -12,6 +12,8 @@ This project "profanity2" was forked from the original project and modified to g Project "profanity2" is not generating key anymore, instead it adjusts user-provided public key until desired vanity address will be discovered. Users provide seed public key in form of 128-symbol hex string with `-z` parameter flag. Resulting private key should be used to be added to seed private key to achieve final private key of the desired vanity address (private keys are just 256-bit numbers). Running "profanity2" can even be outsourced to someone completely unreliable - it is still safe by design. +Note: when upgrading to a new version of profanity2, delete the `cache-opencl.*` files (or pass `--no-cache`) once so the OpenCL program is rebuilt with the new kernel. + ## Getting public key for mandatory `-z` parameter Generate private key and public key via openssl in terminal (remove prefix "04" from public key): @@ -262,9 +264,6 @@ reported), and a very short mask can produce more matches per GPU round than the buffer holds — the program prints a warning with the number of dropped matches if that happens. -Note: if you have run an older version of profanity2 before, delete the `cache-opencl.*` -files (or pass `--no-cache`) once so the OpenCL program is rebuilt with the new kernel. - ### Character classes anywhere (`--zeros`, `--letters`, `--numbers`) Score on the total amount of matching characters anywhere in the address: @@ -334,16 +333,17 @@ zeroth transaction** of the found account instead of the account address itself: ``` ### Benchmarks - Current version -|Model|Clock Speed|Memory Speed|Modified straps|Speed|Time to match eight characters -|:-:|:-:|:-:|:-:|:-:|:-:| -|GTX 1070 OC|1950|4450|NO|179.0 MH/s| ~24s -|GTX 1070|1750|4000|NO|163.0 MH/s| ~26s -|RX 480|1328|2000|YES|120.0 MH/s| ~36s -|RTX 4090|-|-|-|1096 MH/s| ~3s -|Apple Silicon M1
(8-core GPU)|-|-|-|45.0 MH/s| ~97s -|Apple Silicon M1 Max
(32-core GPU)|-|-|-|172.0 MH/s| ~25s -|Apple Silicon M3 Pro
(18-core GPU)|-|-|-|97 MH/s| ~45s -|Apple Silicon M4 Max
(40-core GPU)|-|-|-|350 MH/s| ~12s +|Model|Clock Speed|Memory Speed|Speed|Time to match eight characters +|:-:|:-:|:-:|:-:|:-:| +|GTX 1070|1750|4000|225 MH/s| ~19s +|RTX 4090|2550|10500|1361 MH/s| ~3s +|RX 480|1328|4000|120 MH/s| ~36s +|RX 7900 XTX|2500|10000|592 MH/s| ~7s +|Apple Silicon M1
(8-core GPU)|1278|4266|60 MH/s| ~72s +|Apple Silicon M1 Max
(32-core GPU)|1296|6400|229 MH/s| ~19s +|Apple Silicon M2
(10-core GPU)|1398|6400|75 MH/s| ~57s +|Apple Silicon M3 Pro
(18-core GPU)|1398|6400|129 MH/s| ~33s +|Apple Silicon M4 Max
(40-core GPU)|1800|8533|467 MH/s| ~9s # License diff --git a/SpeedSample.cpp b/SpeedSample.cpp index 8a92d9f..dc584d2 100644 --- a/SpeedSample.cpp +++ b/SpeedSample.cpp @@ -27,8 +27,8 @@ double SpeedSample::getSpeed() const { void SpeedSample::sample(const double V) { const timepoint newTime = now(); - auto delta = std::chrono::duration_cast(newTime - m_lastTime).count(); - m_lSpeeds.push_back((1000 * V) / delta); + auto delta = std::chrono::duration_cast(newTime - m_lastTime).count(); + m_lSpeeds.push_back((1000000.0 * V) / delta); m_lastTime = newTime; if (m_lSpeeds.size() > m_length) { m_lSpeeds.pop_front(); diff --git a/profanity.cl b/profanity.cl index 9e2a5af..3559ecd 100644 --- a/profanity.cl +++ b/profanity.cl @@ -48,7 +48,7 @@ #define bswap32(n) (rotate(n & 0x00FF00FF, 24U)|(rotate(n, 8U) & 0x00FF00FF)) typedef uint mp_word; -typedef struct { +typedef struct __attribute__((aligned(16))) { mp_word d[MP_WORDS]; } mp_number; @@ -558,6 +558,10 @@ __kernel void profanity_inverse(__global const mp_number * const pDeltaX, __glob pInverse[id] = copy1; } +static inline uchar profanity_byte(const uint * const address, const int i) { + return (uchar)(address[i >> 2] >> ((i & 3) << 3)); +} + // This kernel performs en elliptical curve point addition. See: // https://en.wikipedia.org/wiki/Elliptic_curve_point_multiplication#Point_addition // I've made one mathematical optimization by never calculating x_r, @@ -628,15 +632,10 @@ __kernel void profanity_inverse(__global const mp_number * const pDeltaX, __glob // in hopes that using constant storage instead of private storage // will aid speeds. // -// After the above point addition this kernel calculates the public address -// corresponding to the point and stores it in pInverse which is used only -// as interim storage as it won't otherwise be used again this cycle. -// -// One of the scoring kernels will run after this and fetch the address -// from pInverse. -__kernel void profanity_iterate(__global mp_number * const pDeltaX, __global mp_number * const pInverse, __global mp_number * const pPrevLambda) { - const size_t id = get_global_id(0); - +// After the above point addition this calculates the public address +// corresponding to the point and returns it in private memory, where the +// scoring below grades it without a round trip through global memory. +static inline void profanity_iterate(__global mp_number * const pDeltaX, __global const mp_number * const pInverse, __global mp_number * const pPrevLambda, const size_t id, const uchar bContract, uint * const address) { // negativeGx = 0x8641998106234453aa5f9d6a3178f4f8fd640324d231d726a60d7ea3e907e497 mp_number negativeGx = { {0xe907e497, 0xa60d7ea3, 0xd231d726, 0xfd640324, 0x3178f4f8, 0xaa5f9d6a, 0x06234453, 0x86419981 } }; @@ -688,15 +687,36 @@ __kernel void profanity_iterate(__global mp_number * const pDeltaX, __global mp_ sha3_keccakf(&h); - // Save public address hash in pInverse, only used as interim storage until next cycle - pInverse[id].d[0] = h.d[3]; - pInverse[id].d[1] = h.d[4]; - pInverse[id].d[2] = h.d[5]; - pInverse[id].d[3] = h.d[6]; - pInverse[id].d[4] = h.d[7]; + // The address is the low 20 bytes of the hash, words 3 through 7. + address[0] = h.d[3]; + address[1] = h.d[4]; + address[2] = h.d[5]; + address[3] = h.d[6]; + address[4] = h.d[7]; + + if (bContract) { + ethhash c = { { 0 } }; + + // set up keccak(0xd6, 0x94, address, 0x80) + c.b[0] = 0xd6; + c.b[1] = 0x94; + for (int i = 0; i < 20; ++i) { + c.b[i + 2] = profanity_byte(address, i); + } + c.b[22] = 0x80; + + c.b[23] ^= 0x01; // length 23 + sha3_keccakf(&c); + + address[0] = c.d[3]; + address[1] = c.d[4]; + address[2] = c.d[5]; + address[3] = c.d[6]; + address[4] = c.d[7]; + } } -void profanity_result_update(const size_t id, __global const uchar * const hash, __global result * const pResult, const uchar score, const uchar scoreMax) { +void profanity_result_update(const size_t id, const uint * const address, __global result * const pResult, const uchar score, const uchar scoreMax) { if (score && score > scoreMax) { uchar hasResult = atomic_inc(&pResult[score].found); // NOTE: If "too many" results are found it'll wrap around to 0 again and overwrite last result. Only relevant if global worksize exceeds MAX(uint). @@ -705,48 +725,29 @@ void profanity_result_update(const size_t id, __global const uchar * const hash, pResult[score].foundId = id; for (int i = 0; i < 20; ++i) { - pResult[score].foundHash[i] = hash[i]; + pResult[score].foundHash[i] = profanity_byte(address, i); } } } } -__kernel void profanity_transform_contract(__global mp_number * const pInverse) { - const size_t id = get_global_id(0); - __global const uchar * const hash = (__global const uchar *)&pInverse[id].d[0]; +// Prevent the compiler from deleting the keccak behind profanity_iterate +// Scores 1 for address(0), which is unreachable, and 0 on everything else +static inline int profanity_score_fn_benchmark(const uint * const address, __constant const uchar * const data1, __constant const uchar * const data2) { + uint sum = 0; - ethhash h; - for (int i = 0; i < 50; ++i) { - h.d[i] = 0; + for (int i = 0; i < 5; ++i) { + sum |= address[i]; } - // set up keccak(0xd6, 0x94, address, 0x80) - h.b[0] = 214; - h.b[1] = 148; - for (int i = 0; i < 20; i++) { - h.b[i + 2] = hash[i]; - } - h.b[22] = 128; - - h.b[23] ^= 0x01; // length 23 - sha3_keccakf(&h); - pInverse[id].d[0] = h.d[3]; - pInverse[id].d[1] = h.d[4]; - pInverse[id].d[2] = h.d[5]; - pInverse[id].d[3] = h.d[6]; - pInverse[id].d[4] = h.d[7]; -} - -__kernel void profanity_score_benchmark(__global mp_number * const pInverse, __global result * const pResult, __constant const uchar * const data1, __constant const uchar * const data2, const uchar scoreMax) { - const size_t id = get_global_id(0); - __global const uchar * const hash = (__global const uchar *)&pInverse[id].d[0]; - int score = 0; - - profanity_result_update(id, hash, pResult, score, scoreMax); + return sum == 0; } // Reports every hash that matches the given mask (data1) and pattern (data2) // exactly, unlike the scoring kernels which report at most one hash per score. +// It has no score to hand back, so it cannot go through PROFANITY_SCORE_KERNEL, +// but it takes the same arguments so that the host can set them without caring +// which kernel the mode selected. scoreMax is unused. // // pResult[0].found counts the matches found by this launch; matches are // appended at pResult[1..PROFANITY_MAX_SCORE] in arrival order. The host reads @@ -754,14 +755,26 @@ __kernel void profanity_score_benchmark(__global mp_number * const pInverse, __g // check below also guarantees that no two work items ever write the same slot. // Matches beyond the buffer capacity are counted but not stored; the host // reports how many were dropped. -__kernel void profanity_exact_match(__global mp_number * const pInverse, __global result * const pResult, __constant const uchar * const data1, __constant const uchar * const data2, const uchar scoreMax) { +__kernel void profanity_iterate_exact_match( + __global mp_number * const pDeltaX, + __global const mp_number * const pInverse, + __global mp_number * const pPrevLambda, + __global result * const pResult, + __constant const uchar * const data1, + __constant const uchar * const data2, + const uchar scoreMax, + const uchar bContract) { const size_t id = get_global_id(0); - __global const uchar * const hash = (__global const uchar *)&pInverse[id].d[0]; + uint address[5]; + profanity_iterate(pDeltaX, pInverse, pPrevLambda, id, bContract, address); - for (int i = 0; i < 20; ++i) { + __constant const uint * const mask = (__constant const uint *)data1; + __constant const uint * const want = (__constant const uint *)data2; + + for (int i = 0; i < 5; ++i) { // Wildcard positions have zero mask bits in data1 and zero bits in // data2, so they compare equal for any hash byte. - if ((hash[i] & data1[i]) != data2[i]) { + if ((address[i] & mask[i]) != want[i]) { return; } } @@ -771,39 +784,37 @@ __kernel void profanity_exact_match(__global mp_number * const pInverse, __globa pResult[matchIndex + 1].foundId = id; for (int i = 0; i < 20; ++i) { - pResult[matchIndex + 1].foundHash[i] = hash[i]; + pResult[matchIndex + 1].foundHash[i] = profanity_byte(address, i); } } } -__kernel void profanity_score_matching(__global mp_number * const pInverse, __global result * const pResult, __constant const uchar * const data1, __constant const uchar * const data2, const uchar scoreMax) { - const size_t id = get_global_id(0); - __global const uchar * const hash = (__global const uchar *)&pInverse[id].d[0]; +static inline int profanity_score_fn_matching(const uint * const address, __constant const uchar * const data1, __constant const uchar * const data2) { int score = 0; for (int i = 0; i < 20; ++i) { - if (data1[i] > 0 && (hash[i] & data1[i]) == data2[i]) { + if (data1[i] > 0 && (profanity_byte(address, i) & data1[i]) == data2[i]) { ++score; } } - profanity_result_update(id, hash, pResult, score, scoreMax); + return score; } -__kernel void profanity_score_leading(__global mp_number * const pInverse, __global result * const pResult, __constant const uchar * const data1, __constant const uchar * const data2, const uchar scoreMax) { - const size_t id = get_global_id(0); - __global const uchar * const hash = (__global const uchar *)&pInverse[id].d[0]; +static inline int profanity_score_fn_leading(const uint * const address, __constant const uchar * const data1, __constant const uchar * const data2) { int score = 0; for (int i = 0; i < 20; ++i) { - if ((hash[i] & 0xF0) >> 4 == data1[0]) { + const uchar byte = profanity_byte(address, i); + + if ((byte & 0xF0) >> 4 == data1[0]) { ++score; } else { break; } - if ((hash[i] & 0x0F) == data1[0]) { + if ((byte & 0x0F) == data1[0]) { ++score; } else { @@ -811,17 +822,16 @@ __kernel void profanity_score_leading(__global mp_number * const pInverse, __glo } } - profanity_result_update(id, hash, pResult, score, scoreMax); + return score; } -__kernel void profanity_score_range(__global mp_number * const pInverse, __global result * const pResult, __constant const uchar * const data1, __constant const uchar * const data2, const uchar scoreMax) { - const size_t id = get_global_id(0); - __global const uchar * const hash = (__global const uchar *)&pInverse[id].d[0]; +static inline int profanity_score_fn_range(const uint * const address, __constant const uchar * const data1, __constant const uchar * const data2) { int score = 0; for (int i = 0; i < 20; ++i) { - const uchar first = (hash[i] & 0xF0) >> 4; - const uchar second = (hash[i] & 0x0F); + const uchar byte = profanity_byte(address, i); + const uchar first = (byte & 0xF0) >> 4; + const uchar second = (byte & 0x0F); if (first >= data1[0] && first <= data2[0]) { ++score; @@ -832,31 +842,28 @@ __kernel void profanity_score_range(__global mp_number * const pInverse, __globa } } - profanity_result_update(id, hash, pResult, score, scoreMax); + return score; } -__kernel void profanity_score_zerobytes(__global mp_number * const pInverse, __global result * const pResult, __constant const uchar * const data1, __constant const uchar * const data2, const uchar scoreMax) { - const size_t id = get_global_id(0); - __global const uchar * const hash = (__global const uchar *)&pInverse[id].d[0]; +static inline int profanity_score_fn_zerobytes(const uint * const address, __constant const uchar * const data1, __constant const uchar * const data2) { int score = 0; for (int i = 0; i < 20; ++i) { - if (hash[i] == 0) { + if (profanity_byte(address, i) == 0) { score++; } } - profanity_result_update(id, hash, pResult, score, scoreMax); + return score; } -__kernel void profanity_score_leadingrange(__global mp_number * const pInverse, __global result * const pResult, __constant const uchar * const data1, __constant const uchar * const data2, const uchar scoreMax) { - const size_t id = get_global_id(0); - __global const uchar * const hash = (__global const uchar *)&pInverse[id].d[0]; +static inline int profanity_score_fn_leadingrange(const uint * const address, __constant const uchar * const data1, __constant const uchar * const data2) { int score = 0; for (int i = 0; i < 20; ++i) { - const uchar first = (hash[i] & 0xF0) >> 4; - const uchar second = (hash[i] & 0x0F); + const uchar byte = profanity_byte(address, i); + const uchar first = (byte & 0xF0) >> 4; + const uchar second = (byte & 0x0F); if (first >= data1[0] && first <= data2[0]) { ++score; @@ -873,20 +880,21 @@ __kernel void profanity_score_leadingrange(__global mp_number * const pInverse, } } - profanity_result_update(id, hash, pResult, score, scoreMax); + return score; } -__kernel void profanity_score_mirror(__global mp_number * const pInverse, __global result * const pResult, __constant const uchar * const data1, __constant const uchar * const data2, const uchar scoreMax) { - const size_t id = get_global_id(0); - __global const uchar * const hash = (__global const uchar *)&pInverse[id].d[0]; +static inline int profanity_score_fn_mirror(const uint * const address, __constant const uchar * const data1, __constant const uchar * const data2) { int score = 0; for (int i = 0; i < 10; ++i) { - const uchar leftLeft = (hash[9 - i] & 0xF0) >> 4; - const uchar leftRight = (hash[9 - i] & 0x0F); + const uchar left = profanity_byte(address, 9 - i); + const uchar right = profanity_byte(address, 10 + i); - const uchar rightLeft = (hash[10 + i] & 0xF0) >> 4; - const uchar rightRight = (hash[10 + i] & 0x0F); + const uchar leftLeft = (left & 0xF0) >> 4; + const uchar leftRight = (left & 0x0F); + + const uchar rightLeft = (right & 0xF0) >> 4; + const uchar rightRight = (right & 0x0F); if (leftRight != rightLeft) { break; @@ -901,16 +909,16 @@ __kernel void profanity_score_mirror(__global mp_number * const pInverse, __glob ++score; } - profanity_result_update(id, hash, pResult, score, scoreMax); + return score; } -__kernel void profanity_score_doubles(__global mp_number * const pInverse, __global result * const pResult, __constant const uchar * const data1, __constant const uchar * const data2, const uchar scoreMax) { - const size_t id = get_global_id(0); - __global const uchar * const hash = (__global const uchar *)&pInverse[id].d[0]; +static inline int profanity_score_fn_doubles(const uint * const address, __constant const uchar * const data1, __constant const uchar * const data2) { int score = 0; for (int i = 0; i < 20; ++i) { - if ((hash[i] == 0x00) || (hash[i] == 0x11) || (hash[i] == 0x22) || (hash[i] == 0x33) || (hash[i] == 0x44) || (hash[i] == 0x55) || (hash[i] == 0x66) || (hash[i] == 0x77) || (hash[i] == 0x88) || (hash[i] == 0x99) || (hash[i] == 0xAA) || (hash[i] == 0xBB) || (hash[i] == 0xCC) || (hash[i] == 0xDD) || (hash[i] == 0xEE) || (hash[i] == 0xFF)) { + const uchar byte = profanity_byte(address, i); + + if ((((byte >> 4) ^ byte) & 0x0f) == 0) { ++score; } else { @@ -918,5 +926,34 @@ __kernel void profanity_score_doubles(__global mp_number * const pInverse, __glo } } - profanity_result_update(id, hash, pResult, score, scoreMax); -} + return score; +} + +// One kernel per scoring mode, each taking a candidate from the point addition +// through to its score. bContract is uniform across the launch and selects the +// second hash that turns a sender into the contract it deploys at nonce 0. +#define PROFANITY_SCORE_KERNEL(NAME) \ +__kernel void profanity_iterate_score_##NAME( \ + __global mp_number * const pDeltaX, \ + __global const mp_number * const pInverse, \ + __global mp_number * const pPrevLambda, \ + __global result * const pResult, \ + __constant const uchar * const data1, \ + __constant const uchar * const data2, \ + const uchar scoreMax, \ + const uchar bContract) { \ + const size_t id = get_global_id(0); \ + uint address[5]; \ + profanity_iterate(pDeltaX, pInverse, pPrevLambda, id, bContract, address); \ + const int score = profanity_score_fn_##NAME(address, data1, data2); \ + profanity_result_update(id, address, pResult, score, scoreMax); \ +} + +PROFANITY_SCORE_KERNEL(benchmark) +PROFANITY_SCORE_KERNEL(matching) +PROFANITY_SCORE_KERNEL(leading) +PROFANITY_SCORE_KERNEL(range) +PROFANITY_SCORE_KERNEL(zerobytes) +PROFANITY_SCORE_KERNEL(leadingrange) +PROFANITY_SCORE_KERNEL(mirror) +PROFANITY_SCORE_KERNEL(doubles) diff --git a/types.hpp b/types.hpp index 0799448..7a9c180 100644 --- a/types.hpp +++ b/types.hpp @@ -14,7 +14,7 @@ typedef cl_uint mp_word; -typedef struct { +typedef struct alignas(16) { mp_word d[MP_NWORDS]; } mp_number;