diff --git a/src/crypto/wbag-hash.c b/src/crypto/wbag-hash.c index 04a80946c..312e2d79b 100644 --- a/src/crypto/wbag-hash.c +++ b/src/crypto/wbag-hash.c @@ -36,6 +36,14 @@ extern "C" { #define WBAG_SCRATCHPAD (1u * 1024u * 1024u) // 1 MiB per-hash #endif +#ifndef WBAG_SCRATCHPAD_RESIDENT +#define WBAG_SCRATCHPAD_RESIDENT (128u * 1024u * 1024u) // 128 MiB resident per-hash +#endif + +#ifndef WBAG_WINDOW_ALIGN +#define WBAG_WINDOW_ALIGN (1u * 1024u * 1024u) // 1 MiB-aligned window +#endif + #define WBAG_N_ENTRIES ( (WBAG_CACHE_BYTES / WBAG_ENTRY_BYTES) ) #ifndef WBAG_ARGON2_P @@ -52,6 +60,15 @@ _Static_assert(WBAG_ARGON2_P == 2u, "WBAG requires Argon2 p=2"); #endif // Compile-time guards +#if (WBAG_SCRATCHPAD_RESIDENT % 8u) != 0 +# error "WBAG_SCRATCHPAD_RESIDENT must be a multiple of 8" +#endif +#if (WBAG_SCRATCHPAD_RESIDENT < WBAG_SCRATCHPAD) +# error "WBAG_SCRATCHPAD_RESIDENT must be >= WBAG_SCRATCHPAD" +#endif +#if (WBAG_WINDOW_ALIGN & (WBAG_WINDOW_ALIGN - 1u)) +# error "WBAG_WINDOW_ALIGN must be a power of two" +#endif #if (WBAG_ENTRY_BYTES < 32) # error "WBAG_ENTRY_BYTES must be at least 32" #endif @@ -182,23 +199,57 @@ static inline uint64_t wbag_mul128_hi_lo_p(uint64_t a, uint64_t b, uint64_t *lo) } static uint64_t (* volatile wbag_mul128_fn)(uint64_t, uint64_t, uint64_t*) = wbag_mul128_hi_lo_p; +// resident fill +static inline uint64_t wbag_seed_from_cache(const uint8_t* cache, size_t cache_bytes) { + // 8 taps across cache; 8 loads total + const size_t step = cache_bytes / 8u; + uint64_t s = 0x6a09e667f3bcc908ULL ^ ((uint64_t)WBAG_params_signature() << 32); + for (size_t i = 0; i < 8; i++) { + const size_t off = (i * step) & ~(size_t)7; + s ^= WBAG_load64_le(cache + off); + s = WBAG_splitmix64(s); + } + return WBAG_splitmix64(s); +} + +uint64_t WBAG_resident_seed_from_cache(const uint8_t* cache, size_t cache_bytes) { + return wbag_seed_from_cache(cache, cache_bytes); +} + +int WBAG_resident_fill(uint8_t *resident, size_t resident_len, uint64_t seed) +{ + if (!resident) { errno = EINVAL; return -1; } + if (resident_len < (size_t)WBAG_SCRATCHPAD) { errno = EINVAL; return -1; } + if ((resident_len % 8u) != 0u) { errno = EINVAL; return -1; } + + uint64_t q = seed ^ 0x243F6A8885A308D3ULL; + for (size_t i = 0; i < resident_len/8; i++) { + q = WBAG_splitmix64(q); + WBAG_store64_le(resident + 8*i, q); + } + return 0; +} + // 32-bit FNV-1a over the textual description of params uint32_t WBAG_params_signature(void) { - char buf[192]; + char buf[256]; int n = snprintf(buf, sizeof(buf), - "cache=%zu;entry=%zu;entries=%zu;rounds=%u;scratch=%zu;cache_prf=argon2d;ver=0x%02x;t=%u;p=%u", + "cache=%zu;entry=%zu;entries=%zu;rounds=%u;" + "scratch=%zu;resident=%zu;win_align=%zu;mul128=%u;" + "cache_prf=argon2d;ver=0x%02x;t=%u;p=%u", (size_t)WBAG_CACHE_BYTES, (size_t)WBAG_ENTRY_BYTES, (size_t)WBAG_N_ENTRIES, - (unsigned)WBAG_ROUNDS, (size_t)WBAG_SCRATCHPAD, + (unsigned)WBAG_ROUNDS, + (size_t)WBAG_SCRATCHPAD, + (size_t)WBAG_SCRATCHPAD_RESIDENT, + (size_t)WBAG_WINDOW_ALIGN, + (unsigned)WBAG_MUL128_EVERY, (unsigned)ARGON2_VERSION_13, (unsigned)WBAG_ARGON2_T, (unsigned)WBAG_ARGON2_P); if (n < 0) n = 0; if (n >= (int)sizeof(buf)) n = (int)sizeof(buf) - 1; uint32_t h = 2166136261u; - for (int i = 0; i < n; i++) { - h ^= (uint8_t)buf[i]; - h *= 16777619u; - } + for (int i = 0; i < n; i++) { h ^= (uint8_t)buf[i]; h *= 16777619u; } return h; } @@ -229,16 +280,24 @@ int WBAG_build_cache(const uint8_t seed32[32], uint8_t *out_base) return 0; } -// Hash: prehash, rounds, Groestl finalize -int WBAG_hash(const void *bd, size_t bd_len, const uint8_t *cache, uint8_t out32[32]) { - if (!out32) { errno = EINVAL; return -1; } - if (!cache) { errno = EINVAL; return -1; } - if (bd_len > 0 && !bd) { errno = EINVAL; return -1; } +// Hash: prehash, rounds with aligned window within resident, Groestl finalize +int WBAG_hash(const void *bd, size_t bd_len, + const uint8_t *cache, + const uint8_t *resident, size_t resident_len, + uint8_t out32[32]) +{ + if (!out32 || !cache) { errno = EINVAL; return -1; } + if (bd_len && !bd) { errno = EINVAL; return -1; } if (bd_len > (size_t)WBAG_MAX_BD_LEN) { errno = EFBIG; return -1; } - const size_t pre_len = 1 + 32; // 'f' || st_bytes - if (bd_len > SIZE_MAX - pre_len) { errno = EOVERFLOW; return -1; } + // resident requirements + size_t need = (size_t)WBAG_SCRATCHPAD_RESIDENT; + if (need < (size_t)WBAG_SCRATCHPAD) need = (size_t)WBAG_SCRATCHPAD; + if (!resident || resident_len < need || (resident_len % 8u)) { errno = EINVAL; return -1; } + // finalize buffer + const size_t pre_len = 1 + 32; + if (bd_len > SIZE_MAX - pre_len) { errno = EOVERFLOW; return -1; } uint8_t *buf = (uint8_t*)malloc(pre_len + bd_len); if (!buf) { errno = ENOMEM; return -1; } @@ -248,7 +307,7 @@ int WBAG_hash(const void *bd, size_t bd_len, const uint8_t *cache, uint8_t out32 memcpy(buf + 1, bd, bd_len); blake256_hash(st_bytes, buf, (uint64_t)(bd_len + 1)); - // 2) State lanes + PRNG seed + scratchpad + // 2) State + PRNG seed uint64_t st64[4]; st64[0] = WBAG_load64_le(st_bytes + 0); st64[1] = WBAG_load64_le(st_bytes + 8); @@ -261,40 +320,31 @@ int WBAG_hash(const void *bd, size_t bd_len, const uint8_t *cache, uint8_t out32 s ^= st64[2]; s = WBAG_splitmix64(s); s ^= st64[3]; s = WBAG_splitmix64(s); - uint8_t *scratchpad = (uint8_t*)malloc(WBAG_SCRATCHPAD); - if (!scratchpad) { free(buf); errno = ENOMEM; return -1; } - { - uint64_t q = st64[0] ^ st64[1] ^ st64[2] ^ st64[3] ^ 0x243F6A8885A308D3ULL; - for (size_t i = 0; i < WBAG_SCRATCHPAD/8; i++) { - q = WBAG_splitmix64(q); - WBAG_store64_le(scratchpad + 8*i, q); - } - } - const size_t smask = (WBAG_SCRATCHPAD/8u) - 1u; + const size_t smask_touch = (WBAG_SCRATCHPAD/8u) - 1u; + const size_t window_span = need - (size_t)WBAG_SCRATCHPAD; - // 3) Dependent random rounds + 128-bit mul-high dependency chain + // 3) Loop const size_t n_mask = (size_t)WBAG_N_ENTRIES - 1; - - const volatile uint32_t rounds = (uint32_t)WBAG_ROUNDS; + const uint32_t rounds = (uint32_t)WBAG_ROUNDS; for (uint32_t t = 0; t < rounds; t++) { - size_t idx = (size_t)(s & (uint64_t)n_mask); + const size_t idx = (size_t)(s & (uint64_t)n_mask); const uint8_t *entry = cache + idx * (size_t)WBAG_ENTRY_BYTES; - uint64_t w0 = WBAG_load64_le(entry + 0); - uint64_t w1 = WBAG_load64_le(entry + 8); - uint64_t w2 = WBAG_load64_le(entry + 16); - uint64_t w3 = WBAG_load64_le(entry + 24); + const uint64_t w0 = WBAG_load64_le(entry + 0); + const uint64_t w1 = WBAG_load64_le(entry + 8); + const uint64_t w2 = WBAG_load64_le(entry + 16); + const uint64_t w3 = WBAG_load64_le(entry + 24); - unsigned r0 = (unsigned)(s & 63u); - unsigned r1 = (unsigned)((s >> 7) & 63u); - unsigned r2 = (unsigned)((s >> 17) & 63u); - unsigned r3 = (unsigned)((s >> 29) & 63u); + const unsigned r0 = (unsigned)(s & 63u); + const unsigned r1 = (unsigned)((s >> 7) & 63u); + const unsigned r2 = (unsigned)((s >> 17) & 63u); + const unsigned r3 = (unsigned)((s >> 29) & 63u); - uint64_t x0 = WBAG_rotl64(st64[0] ^ w0, r0); - uint64_t x1 = WBAG_rotl64(st64[1] ^ w1, r1); - uint64_t x2 = WBAG_rotl64(st64[2] ^ w2, r2); - uint64_t x3 = WBAG_rotl64(st64[3] ^ w3, r3); + const uint64_t x0 = WBAG_rotl64(st64[0] ^ w0, r0); + const uint64_t x1 = WBAG_rotl64(st64[1] ^ w1, r1); + const uint64_t x2 = WBAG_rotl64(st64[2] ^ w2, r2); + const uint64_t x3 = WBAG_rotl64(st64[3] ^ w3, r3); st64[0] = x0 + x2; st64[1] = x1 + x3; @@ -304,25 +354,30 @@ int WBAG_hash(const void *bd, size_t bd_len, const uint8_t *cache, uint8_t out32 uint64_t mh_mix = 0; if ((t & (WBAG_MUL128_EVERY - 1u)) == 0u) { uint64_t lo128; - uint64_t hi128 = wbag_mul128_fn(st64[0] ^ w2, st64[3] ^ w1, &lo128); - + const uint64_t hi128 = wbag_mul128_fn(st64[0] ^ w2, st64[3] ^ w1, &lo128); st64[2] ^= hi128; - st64[3] = WBAG_rotl64(st64[3] + lo128, (unsigned)(hi128 & 63u)); - + st64[3] = WBAG_rotl64(st64[3] + lo128, (unsigned)(hi128 & 63u)); s ^= WBAG_rotl64(hi128 ^ lo128, (unsigned)((s >> 7) & 63u)); mh_mix = hi128 ^ WBAG_rotl64(lo128, 13); } - size_t sidx = (size_t)(((s >> 11) ^ mh_mix) & (uint64_t)smask); - uint64_t v = WBAG_load64_le(scratchpad + (sidx * 8u)); + size_t window_base = 0; + if (window_span) { + const size_t align = (size_t)WBAG_WINDOW_ALIGN; + const size_t max_base = window_span & ~(align - 1u); + uint64_t rnd = WBAG_rotl64(s ^ mh_mix, 17) ^ WBAG_rotl64(s + (uint64_t)t, 41); + window_base = (size_t)(rnd % (max_base + 1u)); + window_base &= ~(align - 1u); + } + + const size_t sidx = (size_t)(((s >> 11) ^ mh_mix) & (uint64_t)smask_touch); + const uint64_t v = WBAG_load64_le(resident + window_base + (sidx * 8u)); st64[0] ^= v; s ^= WBAG_rotl64(v, (unsigned)((s >> 23) & 63u)); - size_t off = (size_t)(((s >> 3) ^ WBAG_rotl64(mh_mix, 7)) & (uint64_t)((size_t)WBAG_ENTRY_BYTES - 8)); - uint64_t e = WBAG_load64_le(entry + off); - - s ^= e; - s = WBAG_splitmix64(s); + const size_t off = (size_t)(((s >> 3) ^ WBAG_rotl64(mh_mix, 7)) & (uint64_t)((size_t)WBAG_ENTRY_BYTES - 8)); + const uint64_t e = WBAG_load64_le(entry + off); + s = WBAG_splitmix64(s ^ e); } // 4) Finalize: Groestl256('f' || st || bd) @@ -336,19 +391,20 @@ int WBAG_hash(const void *bd, size_t bd_len, const uint8_t *cache, uint8_t out32 memcpy(buf + (1 + 32), bd, bd_len); groestl(buf, (1 + 32) + bd_len, out32); - free(scratchpad); free(buf); return 0; } // Public parameter accessors -size_t WBAG_cache_bytes(void) { return (size_t)WBAG_CACHE_BYTES; } -size_t WBAG_entry_bytes(void) { return (size_t)WBAG_ENTRY_BYTES; } -size_t WBAG_cache_entries(void) { return (size_t)WBAG_N_ENTRIES; } -size_t WBAG_rounds(void) { return (size_t)WBAG_ROUNDS; } -size_t WBAG_scratchpad(void) { return (size_t)WBAG_SCRATCHPAD; } -size_t WBAG_argon_p(void) { return (size_t)WBAG_ARGON2_P; } -size_t WBAG_argon_t(void) { return (size_t)WBAG_ARGON2_T; } +size_t WBAG_cache_bytes(void) { return (size_t)WBAG_CACHE_BYTES; } +size_t WBAG_entry_bytes(void) { return (size_t)WBAG_ENTRY_BYTES; } +size_t WBAG_cache_entries(void) { return (size_t)WBAG_N_ENTRIES; } +size_t WBAG_rounds(void) { return (size_t)WBAG_ROUNDS; } +size_t WBAG_scratchpad_resident(void) { return (size_t)WBAG_SCRATCHPAD_RESIDENT; } +size_t WBAG_window_align(void) { return (size_t)WBAG_WINDOW_ALIGN; } +size_t WBAG_scratchpad(void) { return (size_t)WBAG_SCRATCHPAD; } +size_t WBAG_argon_p(void) { return (size_t)WBAG_ARGON2_P; } +size_t WBAG_argon_t(void) { return (size_t)WBAG_ARGON2_T; } #ifdef __cplusplus } // extern "C" diff --git a/src/crypto/wbag-hash.h b/src/crypto/wbag-hash.h index 9a132d99a..1d096e98f 100644 --- a/src/crypto/wbag-hash.h +++ b/src/crypto/wbag-hash.h @@ -12,7 +12,11 @@ void WBAG_epoch_seed(uint64_t epoch, const uint8_t anchor_hash[32], uint8_t out3 int WBAG_build_cache(const uint8_t seed32[32], uint8_t *out_base); -int WBAG_hash(const void *bd, size_t bd_len, const uint8_t *cache, uint8_t out32[32]); +int WBAG_resident_fill(uint8_t *resident, size_t resident_len, uint64_t seed); + +int WBAG_hash(const void *bd, size_t bd_len, const uint8_t *cache, const uint8_t *resident, size_t resident_len, uint8_t out32[32]); + +uint64_t WBAG_resident_seed_from_cache(const uint8_t* cache, size_t cache_bytes); uint32_t WBAG_params_signature(void); @@ -20,6 +24,8 @@ size_t WBAG_cache_bytes(void); size_t WBAG_entry_bytes(void); size_t WBAG_cache_entries(void); size_t WBAG_rounds(void); +size_t WBAG_scratchpad_resident(void); +size_t WBAG_window_align(void); size_t WBAG_scratchpad(void); size_t WBAG_argon_p(void); size_t WBAG_argon_t(void); diff --git a/tests/hash/CMakeLists.txt b/tests/hash/CMakeLists.txt index 95a6d3957..c55092aeb 100644 --- a/tests/hash/CMakeLists.txt +++ b/tests/hash/CMakeLists.txt @@ -79,6 +79,3 @@ add_test( NAME "hash-variant2-int-sqrt" COMMAND hash-tests "variant2_int_sqrt") -add_test( - NAME "hash-wbag-selftest" - COMMAND hash-tests "wbag_test") diff --git a/tests/hash/main.cpp b/tests/hash/main.cpp index 8d112c924..f4142d1da 100644 --- a/tests/hash/main.cpp +++ b/tests/hash/main.cpp @@ -135,9 +135,6 @@ int main(int argc, char *argv[]) { return 0; } - if (argc == 2 && strcmp(argv[1], "--wbag-full") == 0) - return wbag_test(); - if (argc >= 2 && strcmp(argv[1], "--wbag-hr") == 0) return wbag_hashrate_cli(argc, argv); diff --git a/tests/hash/wbag_gpu_bench.cu b/tests/hash/wbag_gpu_bench.cu index 7d7c59954..42a77fdf9 100644 --- a/tests/hash/wbag_gpu_bench.cu +++ b/tests/hash/wbag_gpu_bench.cu @@ -1,197 +1,272 @@ #include #include #include -#include -#include #include +#include +#include + #include "crypto/wbag-hash.h" #include "cryptonote_basic/wbag_cache.h" -//device utils -__device__ __forceinline__ uint64_t rotl64(uint64_t x, unsigned r){ return (x << (r & 63)) | (x >> ((64 - r) & 63)); } -__device__ __forceinline__ uint64_t splitmix64(uint64_t x){ +// ---------------- device helpers ---------------- +__device__ __forceinline__ uint64_t d_rotl64(uint64_t x, unsigned r) { + return (x << (r & 63)) | (x >> ((64 - r) & 63)); +} +__device__ __forceinline__ uint64_t d_splitmix64(uint64_t x) { x += 0x9e3779b97f4a7c15ULL; x = (x ^ (x >> 30)) * 0xbf58476d1ce4e5b9ULL; x = (x ^ (x >> 27)) * 0x94d049bb133111ebULL; return x ^ (x >> 31); } -__device__ __forceinline__ uint64_t mulhi64(uint64_t a, uint64_t b){ return __umul64hi(a,b); } -__device__ __forceinline__ uint64_t load64_le(const uint8_t* p){ - uint64_t v; memcpy(&v, p, 8); return v; // device memcpy is fine for POD +__device__ __forceinline__ uint64_t d_mulhi64(uint64_t a, uint64_t b) { + return __umul64hi(a, b); +} +__device__ __forceinline__ uint64_t d_load64_le(const uint8_t* p) { + uint64_t v; memcpy(&v, p, 8); return v; +} +__device__ __forceinline__ void d_store64_le(uint8_t* p, uint64_t x) { + memcpy(p, &x, 8); } -__device__ __forceinline__ void store64_le(uint8_t* p, uint64_t x){ memcpy(p, &x, 8); } -// kernel: approximates WBAG_hash hot loop; good for hashrate; not for correctness vectors +// ---------------- kernels ---------------- + +// Per-thread resident fill: same as WBAG_resident_fill +extern "C" __global__ +void wbag_fill_resident(uint8_t* __restrict__ resident, + size_t resident_len, + uint64_t seed, + int nthreads) +{ + const int tid = blockIdx.x * blockDim.x + threadIdx.x; + if (tid >= nthreads) return; + uint8_t* base = resident + (size_t)tid * resident_len; + + uint64_t q = seed ^ 0x243F6A8885A308D3ULL; + for (size_t i = 0; i < resident_len; i += 8) { + q = d_splitmix64(q); + d_store64_le(base + i, q); + } +} + +// Hashrate kernel: WBAG hot loop with aligned window over per-thread resident extern "C" __global__ void wbag_kernel(const uint8_t* __restrict__ cache, size_t entry_bytes, size_t n_entries, uint32_t rounds, uint32_t mul128_every, - uint8_t* __restrict__ scratch_base, size_t scratch_stride, - uint64_t seed, int trials, int nthreads, + const uint8_t* __restrict__ resident, size_t resident_len, + size_t scratchpad_len, size_t window_align, + uint64_t base_seed, int trials, int nthreads, uint64_t* __restrict__ out_cksum) { const int tid = blockIdx.x * blockDim.x + threadIdx.x; if (tid >= nthreads) return; - uint8_t* scratch = scratch_base + size_t(tid) * scratch_stride; - const size_t smask = (scratch_stride/8) - 1; - const size_t n_mask = n_entries - 1; + + const uint8_t* my_resident = resident + (size_t)tid * resident_len; + const size_t n_mask = n_entries - 1; + const size_t smask_touch = (scratchpad_len / 8) - 1; + const size_t window_span = (resident_len > scratchpad_len) + ? (resident_len - scratchpad_len) : 0; + + // Fake prehash; good enough for HR + uint64_t s = d_splitmix64(base_seed ^ (uint64_t)tid); + uint64_t st0 = d_splitmix64(s); + uint64_t st1 = d_splitmix64(st0); + uint64_t st2 = d_splitmix64(st1); + uint64_t st3 = d_splitmix64(st2); + uint64_t checksum = 0; - // per-thread init - uint64_t s = splitmix64(seed ^ uint64_t(tid) ^ 0x243F6A8885A308D3ULL); - uint64_t st0 = splitmix64(s), st1 = splitmix64(st0), st2 = splitmix64(st1), st3 = splitmix64(st2); - - // fill scratchpad sequentially - uint64_t q = st0 ^ st1 ^ st2 ^ st3 ^ 0xDEADBEEFCAFEBABEULL; - for (size_t i = 0; i < scratch_stride; i += 8) { - q = splitmix64(q); - store64_le(scratch + i, q); - } - - // trials for (int ttry = 0; ttry < trials; ++ttry) { - // fresh s each trial to avoid repeating addresses - uint64_t ss = splitmix64(s ^ (uint64_t)ttry); - uint64_t a0=st0, a1=st1, a2=st2, a3=st3; + uint64_t ss = d_splitmix64(s ^ (uint64_t)ttry); + uint64_t a0 = st0, a1 = st1, a2 = st2, a3 = st3; for (uint32_t t = 0; t < rounds; ++t) { - const size_t idx = size_t(ss) & n_mask; + const size_t idx = (size_t)ss & n_mask; const uint8_t* entry = cache + idx * entry_bytes; - uint64_t w0 = load64_le(entry + 0); - uint64_t w1 = load64_le(entry + 8); - uint64_t w2 = load64_le(entry + 16); - uint64_t w3 = load64_le(entry + 24); - unsigned r0 = unsigned(ss & 63u); - unsigned r1 = unsigned((ss >> 7) & 63u); - unsigned r2 = unsigned((ss >> 17) & 63u); - unsigned r3 = unsigned((ss >> 29) & 63u); + const uint64_t w0 = d_load64_le(entry + 0); + const uint64_t w1 = d_load64_le(entry + 8); + const uint64_t w2 = d_load64_le(entry + 16); + const uint64_t w3 = d_load64_le(entry + 24); - uint64_t x0 = rotl64(a0 ^ w0, r0); - uint64_t x1 = rotl64(a1 ^ w1, r1); - uint64_t x2 = rotl64(a2 ^ w2, r2); - uint64_t x3 = rotl64(a3 ^ w3, r3); + const unsigned r0 = unsigned(ss & 63u); + const unsigned r1 = unsigned((ss >> 7) & 63u); + const unsigned r2 = unsigned((ss >> 17) & 63u); + const unsigned r3 = unsigned((ss >> 29) & 63u); + + const uint64_t x0 = d_rotl64(a0 ^ w0, r0); + const uint64_t x1 = d_rotl64(a1 ^ w1, r1); + const uint64_t x2 = d_rotl64(a2 ^ w2, r2); + const uint64_t x3 = d_rotl64(a3 ^ w3, r3); a0 = x0 + x2; a1 = x1 + x3; - a2 = rotl64(x2 ^ (x0 + 0x9e3779b97f4a7c15ULL), 17); - a3 = rotl64(x3 ^ (x1 + 0x94d049bb133111ebULL), 29); + a2 = d_rotl64(x2 ^ (x0 + 0x9e3779b97f4a7c15ULL), 17); + a3 = d_rotl64(x3 ^ (x1 + 0x94d049bb133111ebULL), 29); uint64_t mh_mix = 0; if ((t & (mul128_every - 1u)) == 0u) { - uint64_t lo = a0 ^ w2; // keep “lo” alive - uint64_t hi = mulhi64(a0 ^ w2, a3 ^ w1); + const uint64_t A = (a0 ^ w2); + const uint64_t B = (a3 ^ w1); + const uint64_t lo = A * B; + const uint64_t hi = d_mulhi64(A, B); a2 ^= hi; - a3 = rotl64(a3 + lo, unsigned(hi & 63u)); - ss ^= rotl64(hi ^ lo, unsigned((ss >> 7) & 63u)); - mh_mix = hi ^ rotl64(lo, 13); + a3 = d_rotl64(a3 + lo, unsigned(hi & 63u)); + ss ^= d_rotl64(hi ^ lo, unsigned((ss >> 7) & 63u)); + mh_mix = hi ^ d_rotl64(lo, 13); } - const size_t sidx = size_t(((ss >> 11) ^ mh_mix) & smask); - const uint64_t v = load64_le(scratch + (sidx * 8)); - a0 ^= v; - ss ^= rotl64(v, unsigned((ss >> 23) & 63u)); + size_t window_base = 0; + if (window_span) { + const size_t align = window_align; + const size_t max_base = window_span & ~(align - 1u); + const uint64_t rnd = d_rotl64(ss ^ mh_mix, 17) ^ d_rotl64(ss + (uint64_t)t, 41); + window_base = (size_t)(rnd % (max_base + 1u)); + window_base &= ~(align - 1u); + } - const size_t off = size_t(((ss >> 3) ^ rotl64(mh_mix, 7)) & (entry_bytes - 8)); - const uint64_t e = load64_le(entry + off); - ss = splitmix64(ss ^ e); + const size_t sidx = (size_t)(((ss >> 11) ^ mh_mix) & (uint64_t)smask_touch); + const uint64_t v = d_load64_le(my_resident + window_base + (sidx * 8)); + a0 ^= v; + ss ^= d_rotl64(v, unsigned((ss >> 23) & 63u)); + + const size_t off = (size_t)(((ss >> 3) ^ d_rotl64(mh_mix, 7)) & (uint64_t)(entry_bytes - 8)); + const uint64_t e = d_load64_le(entry + off); + ss = d_splitmix64(ss ^ e); } - // cheap final fold to keep work observable checksum ^= (a0 + (a1 << 1) + (a2 << 2) + (a3 << 3)); } out_cksum[tid] = checksum; } -// host wrapper +// ---------------- host wrapper ---------------- extern "C" int wbag_gpu_bench_once(const cryptonote::wbag::CacheView& view, int threads, int trials, double& hashes_per_sec_out) { + if (threads <= 0 || trials <= 0) return 1; + + const size_t entry_bytes = WBAG_entry_bytes(); + const size_t cache_bytes = view.size(); + const size_t n_entries = cache_bytes / entry_bytes; + const uint32_t rounds = (uint32_t)WBAG_rounds(); + const uint32_t mul128_every = (uint32_t)2u; // WBAG_MUL128_EVERY default + const size_t scratchpad_len = WBAG_scratchpad(); + const size_t resident_len = WBAG_scratchpad_resident(); // 128 MiB by default + const size_t window_align = WBAG_window_align(); + + // Cache upload (persist across calls) static uint8_t* d_cache = nullptr; static size_t d_cache_size = 0; - uint8_t* d_scratch = nullptr; - uint64_t* d_out = nullptr; - - const size_t cache_bytes = view.size(); - const size_t entry_bytes = WBAG_entry_bytes(); - const size_t n_entries = cache_bytes / entry_bytes; - const uint32_t rounds = (uint32_t)WBAG_rounds(); - const uint32_t mul128_every = (uint32_t)2u; // matches default; pass-through if you expose it - const size_t scratch_stride = WBAG_scratchpad(); - - // upload cache once if (!d_cache || d_cache_size != cache_bytes) { - if (d_cache) cudaFree(d_cache); - cudaMalloc(&d_cache, cache_bytes); - cudaMemcpy(d_cache, view.data(), cache_bytes, cudaMemcpyHostToDevice); + if (d_cache) { cudaFree(d_cache); d_cache = nullptr; d_cache_size = 0; } + if (cudaMalloc(&d_cache, cache_bytes) != cudaSuccess) { + std::fprintf(stderr, "GPU: cache alloc failed (%zu bytes)\n", cache_bytes); + return 2; + } + if (cudaMemcpy(d_cache, view.data(), cache_bytes, cudaMemcpyHostToDevice) != cudaSuccess) { + std::fprintf(stderr, "GPU: cache memcpy failed\n"); + cudaFree(d_cache); d_cache = nullptr; d_cache_size = 0; + return 3; + } d_cache_size = cache_bytes; } - // check VRAM for scratchpads - const size_t scratch_total = scratch_stride * size_t(threads); - if (cudaMalloc(&d_scratch, scratch_total) != cudaSuccess) { - fprintf(stderr, "GPU scratch alloc failed (%zu bytes). Reduce --threads.\n", scratch_total); - return 3; - } - cudaMalloc(&d_out, sizeof(uint64_t) * size_t(threads)); - - // launch - const int block = 256; - const int grid = (threads + block - 1) / block; - uint64_t seed = 0xBADC0FFEEULL; - - cudaEvent_t ev0, ev1; cudaEventCreate(&ev0); cudaEventCreate(&ev1); - cudaEventRecord(ev0); - - wbag_kernel<<>>(d_cache, entry_bytes, n_entries, rounds, mul128_every, - d_scratch, scratch_stride, seed, trials, threads, d_out); - - // catch launch errors early - if (auto e = cudaGetLastError(); e != cudaSuccess) { - fprintf(stderr, "GPU kernel launch failed: %s\n", cudaGetErrorString(e)); - cudaFree(d_out); cudaFree(d_scratch); - cudaEventDestroy(ev0); cudaEventDestroy(ev1); + // Per-thread resident allocation + const size_t resident_total = resident_len * (size_t)threads; + size_t free_b = 0, total_b = 0; + cudaMemGetInfo(&free_b, &total_b); + if (resident_total + (64ULL<<20) > free_b) { // keep some headroom + std::fprintf(stderr, "GPU: not enough VRAM for %d threads x %zu MiB (need ~%zu MiB, free %zu MiB)\n", + threads, resident_len >> 20, resident_total >> 20, free_b >> 20); return 4; } - // catch runtime errors (OOB, etc.) - if (auto e = cudaDeviceSynchronize(); e != cudaSuccess) { - fprintf(stderr, "GPU kernel runtime error: %s\n", cudaGetErrorString(e)); - cudaFree(d_out); cudaFree(d_scratch); - cudaEventDestroy(ev0); cudaEventDestroy(ev1); + uint8_t* d_resident = nullptr; + if (cudaMalloc(&d_resident, resident_total) != cudaSuccess) { + std::fprintf(stderr, "GPU: resident alloc failed (%zu bytes)\n", resident_total); return 5; } + // Seed derived from real cache + const uint64_t base_seed = WBAG_resident_seed_from_cache(view.data(), view.size()); + + // Fill per-thread resident (not timed) + { + const int block = 256; + const int grid = (threads + block - 1) / block; + wbag_fill_resident<<>>(d_resident, resident_len, base_seed, threads); + if (auto e = cudaGetLastError(); e != cudaSuccess) { + std::fprintf(stderr, "GPU: fill kernel launch failed: %s\n", cudaGetErrorString(e)); + cudaFree(d_resident); + return 6; + } + if (auto e = cudaDeviceSynchronize(); e != cudaSuccess) { + std::fprintf(stderr, "GPU: fill kernel runtime error: %s\n", cudaGetErrorString(e)); + cudaFree(d_resident); + return 7; + } + } + + // Output buffer + uint64_t* d_out = nullptr; + if (cudaMalloc(&d_out, sizeof(uint64_t) * (size_t)threads) != cudaSuccess) { + std::fprintf(stderr, "GPU: out alloc failed\n"); + cudaFree(d_resident); + return 8; + } + + // Time only the hashing kernel + cudaEvent_t ev0, ev1; cudaEventCreate(&ev0); cudaEventCreate(&ev1); + cudaEventRecord(ev0); + + const int block = 256; + const int grid = (threads + block - 1) / block; + + wbag_kernel<<>>(d_cache, entry_bytes, n_entries, rounds, mul128_every, + d_resident, resident_len, + scratchpad_len, window_align, + base_seed, trials, threads, d_out); + + if (auto e = cudaGetLastError(); e != cudaSuccess) { + std::fprintf(stderr, "GPU: hash kernel launch failed: %s\n", cudaGetErrorString(e)); + cudaFree(d_out); cudaFree(d_resident); + cudaEventDestroy(ev0); cudaEventDestroy(ev1); + return 9; + } + if (auto e = cudaDeviceSynchronize(); e != cudaSuccess) { + std::fprintf(stderr, "GPU: hash kernel runtime error: %s\n", cudaGetErrorString(e)); + cudaFree(d_out); cudaFree(d_resident); + cudaEventDestroy(ev0); cudaEventDestroy(ev1); + return 10; + } + cudaEventRecord(ev1); cudaEventSynchronize(ev1); float ms = 0.f; cudaEventElapsedTime(&ms, ev0, ev1); - std::vector h_out; - h_out.resize(static_cast(threads)); - if (auto e = cudaMemcpy(h_out.data(), d_out, - sizeof(uint64_t) * static_cast(threads), - cudaMemcpyDeviceToHost); - e != cudaSuccess) { - fprintf(stderr, "GPU memcpy failed: %s\n", cudaGetErrorString(e)); - cudaFree(d_out); cudaFree(d_scratch); + // Sanity + std::vector h_out((size_t)threads); + if (cudaMemcpy(h_out.data(), d_out, sizeof(uint64_t) * (size_t)threads, cudaMemcpyDeviceToHost) != cudaSuccess) { + std::fprintf(stderr, "GPU: out memcpy failed\n"); + cudaFree(d_out); cudaFree(d_resident); cudaEventDestroy(ev0); cudaEventDestroy(ev1); - return 6; + return 11; } - - uint64_t acc = 0; for (const uint64_t v : h_out) acc |= v; - if (acc == 0) { - fprintf(stderr, "GPU checksum is zero. Investigate kernel path.\n"); - cudaFree(d_out); cudaFree(d_scratch); + uint64_t acc = 0; for (uint64_t v : h_out) acc |= v; + if (!acc) { + std::fprintf(stderr, "GPU: zero checksum\n"); + cudaFree(d_out); cudaFree(d_resident); cudaEventDestroy(ev0); cudaEventDestroy(ev1); - return 7; + return 12; } const double sec = std::max(1e-6, double(ms) / 1000.0); hashes_per_sec_out = (double(trials) * double(threads)) / sec; cudaFree(d_out); - cudaFree(d_scratch); + cudaFree(d_resident); cudaEventDestroy(ev0); cudaEventDestroy(ev1); return 0; } diff --git a/tests/hash/wbag_test.cpp b/tests/hash/wbag_test.cpp index dd7247b4a..e5404e01a 100644 --- a/tests/hash/wbag_test.cpp +++ b/tests/hash/wbag_test.cpp @@ -1,631 +1,163 @@ #include -#include +#include #include -#include -#include +#include +#include #include -#include #include -#include +#include +#include #include #include -#include +#include #include #include -#include -#include -#include -#include -#include -#include -#if defined(_MSC_VER) && !defined(__clang__) - #include - #pragma intrinsic(__popcnt) -#endif - -static inline unsigned pc8(unsigned x) { -#if defined(_MSC_VER) && !defined(__clang__) - return __popcnt(x); -#elif defined(__GNUC__) || defined(__clang__) - return __builtin_popcount(x); -#else - // SW fallback - x = x - ((x >> 1) & 0x55555555u); - x = (x & 0x33333333u) + ((x >> 2) & 0x33333333u); - return (((x + (x >> 4)) & 0x0F0F0F0Fu) * 0x01010101u) >> 24; -#endif -} - -#include "misc_log_ex.h" -#include "warnings.h" -#include "string_tools.h" -#include "common/util.h" -#include "crypto/hash.h" -#include "crypto/wbag-hash.h" -#include "cryptonote_basic/wbag_cache.h" #include "wbag_test.h" - -#if defined(WBAG_HAVE_WBAG_GPU_BENCH) -extern "C" int wbag_gpu_bench_once(const cryptonote::wbag::CacheView& view, int threads, int trials, double& hashes_per_sec_out); -#endif +#include "crypto/hash.h" +#include "string_tools.h" +#include "cryptonote_basic/wbag_cache.h" extern "C" { + #include "crypto/wbag-hash.h" } -using namespace std; -using namespace crypto; - -// Hardcoded data -struct wbag_anchor_sample { - uint64_t height; - const char *block_id_hex; -}; - -static constexpr wbag_anchor_sample k_wbag_anchor_samples[] = { +// ---------------------- Anchors ---------------------- +const WbagAnchorSample k_wbag_anchor_samples[] = { {770976, "731cabafc8626ce567be7d2ce161d0fefdaa0e6a4f88275513f199362b8d53fd"}, {770832, "c8710d1ff2b98134c4777e1bbc8dc90ddb6d55a2be850dc904c20218cdb5a970"}, {770688, "f54b797ed2c44e2c37a455118df2eef4656d159ac18ed8aa5de1affac2f5c20b"}, {770544, "5197fb807e62e83a5e50e155ddb9975066aca574e22aa6e2f5ccc0a380f1618a"}, {770400, "4b2981fd2b65aeeb28181ba37c1e49181db4123060f1fa205124aa78e77e4963"}, }; +const int k_wbag_anchor_count = + int(sizeof(k_wbag_anchor_samples) / sizeof(k_wbag_anchor_samples[0])); -static const std::unordered_map kGoldenCacheSHA256 = { - /*{770976, "65cbcf2076cf8aaaa221ae6cd3546e3a2e8d1fa125e8af0ab1ca09587685003d"}, - {770832, "4628b4c3ccb4f89ec8298df4aeced02602f31b3b038b6177c751b255c7e40568"}, - {770688, "c44b615591e420ee63bf96a7cb5f87164f6b73d8d5980297b8fe0a76fb4eea4d"}, - {770544, "9a724116ae70501b1b22b0ff517ce5293c9a18547aa0199694c4358c17c226a9"}, - {770400, "24e47c25df4a8d94a826326b8048aebca69107399016b597a74bf4275d775d0e"},*/ -}; - -// Helpers -static inline bool is_hex_digit(char c) { - return (c >= '0' && c <= '9') || (c >= 'a' && c <= 'f') || (c >= 'A' && c <= 'F'); -} - -static std::string trim(const std::string& s) { - const auto a = s.find_first_not_of(" \t\r\n"); - if (a == std::string::npos) return {}; - const auto b = s.find_last_not_of(" \t\r\n"); - return s.substr(a, b - a + 1); -} - -static std::string lower_hex(const uint8_t* p, size_t n) { +// ---------------------- Helpers ---------------------- +static inline std::string print_rate(double hs) { std::ostringstream o; - o << std::hex << std::setfill('0'); - for (size_t i = 0; i < n; ++i) o << std::setw(2) << unsigned(p[i]); + if (hs >= 1000.0) o << std::fixed << std::setprecision(2) << (hs / 1000.0) << " kH/s"; + else o << std::fixed << std::setprecision(0) << hs << " H/s"; return o.str(); } -static inline int hex_nibble(char c) { - if (c >= '0' && c <= '9') return c - '0'; - if (c >= 'a' && c <= 'f') return c - 'a' + 10; - if (c >= 'A' && c <= 'F') return c - 'A' + 10; - return -1; -} - -static std::vector hex_to_bytes(const std::string& hex) { - if ((hex.size() & 1) != 0) { - std::cerr << "vectors: odd-length hex input\n"; std::exit(1); - } - std::vector out(hex.size()/2); - for (size_t i = 0, j = 0; i < hex.size(); i += 2, ++j) { - const int hi = hex_nibble(hex[i]); - const int lo = hex_nibble(hex[i+1]); - if (hi < 0 || lo < 0) { std::cerr << "vectors: non-hex char in input\n"; std::exit(1); } - out[j] = static_cast((hi << 4) | lo); - } - return out; -} - -static bool ensure_wbag_cache() +// ---------------------- Cache setup via real code ---------------------- +static bool ensure_wbag_cache(int anchor_index, + WbagCacheView& out_cpu_view, + cryptonote::wbag::CacheView& out_gpu_view, + double& gen_seconds) { - static std::once_flag once; - static bool ready = false; + using cryptonote::wbag::ensure_cache; + using cryptonote::wbag::cache_view; + using cryptonote::wbag::epoch; - std::call_once(once, []{ - for (const auto &sample : k_wbag_anchor_samples) { - crypto::hash anchor{}; - if (!epee::string_tools::hex_to_pod(sample.block_id_hex, anchor)) continue; - const uint64_t ep = cryptonote::wbag::epoch(sample.height); - if (!cryptonote::wbag::ensure_cache(ep, anchor)) continue; - auto view = cryptonote::wbag::cache_view(); - if (view && view.data()) { ready = true; break; } - } - }); + if (anchor_index < 0 || anchor_index >= k_wbag_anchor_count) return false; - auto view = cryptonote::wbag::cache_view(); - return ready && view && view.data(); -} + const auto& a = k_wbag_anchor_samples[anchor_index]; + crypto::hash anchor{}; + if (!epee::string_tools::hex_to_pod(a.block_id_hex, anchor)) return false; -void wbag_hash(const void *data, size_t length, char *hash) -{ - if (!ensure_wbag_cache()) - throw std::runtime_error("WBAG cache unavailable"); + const uint64_t ep = epoch(a.height); - auto cached = cryptonote::wbag::cache_view(); - if (!cached || !cached.data()) - throw std::runtime_error("WBAG cache view invalid"); + const auto t0 = std::chrono::high_resolution_clock::now(); + if (!ensure_cache(ep, anchor)) return false; + gen_seconds = std::chrono::duration(std::chrono::high_resolution_clock::now() - t0).count(); - if (WBAG_hash(data, length, cached.data(), reinterpret_cast(hash)) != 0) - throw std::runtime_error("WBAG_hash failed"); -} + out_gpu_view = cache_view(); + if (!out_gpu_view || !out_gpu_view.data() || out_gpu_view.size() == 0) return false; -static bool write_cache_and_sha256(const cryptonote::wbag::CacheView& view, std::string& sha_hex, std::string& note) { - static constexpr const char *k_cache_file = "wbag_cache.dat"; - { - std::ofstream out(k_cache_file, std::ios::binary | std::ios::trunc); - if (!out) { note = "write cache file failed"; return false; } - out.write(reinterpret_cast(view.data()), static_cast(view.size())); - if (!out.good()) { note = "write cache bytes failed"; return false; } - } + const size_t resident_bytes = WBAG_scratchpad_resident(); + static std::vector resident; + resident.resize(resident_bytes); - crypto::hash file_hash{}; - if (!tools::sha256sum(k_cache_file, file_hash)) { - note = "sha256sum failed"; - (void)std::remove(k_cache_file); - return false; - } - sha_hex = epee::string_tools::pod_to_hex(file_hash); - (void)std::remove(k_cache_file); + const uint64_t rseed = WBAG_resident_seed_from_cache( + out_gpu_view.data(), out_gpu_view.size()); + if (WBAG_resident_fill(resident.data(), resident.size(), rseed) != 0) return false; + + out_cpu_view.cache = out_gpu_view.data(); + out_cpu_view.cache_size = out_gpu_view.size(); + out_cpu_view.resident = resident.data(); + out_cpu_view.resident_size = resident.size(); return true; } -static void benchmark_hash(const cryptonote::wbag::CacheView& view, int trials, const std::vector& block, double& perf_khs_out, std::string& last_digest_hex_out) { - std::vector out(crypto::HASH_SIZE); - - // Warm-up - for (int i = 0; i < 64; ++i) - WBAG_hash(block.data(), block.size(), view.data(), out.data()); - - // Timed runs: median of 3 - double times[3]; - for (int r = 0; r < 3; ++r) { - const auto t0 = std::chrono::high_resolution_clock::now(); - if (trials <= 512) { - for (int i = 0; i < trials; ++i) { - WBAG_hash(block.data(), block.size(), view.data(), out.data()); - if (((i + 1) % 8) == 0) std::cerr << "." << std::flush; - } - } else { - for (int i = 0; i < trials; ++i) - WBAG_hash(block.data(), block.size(), view.data(), out.data()); - } - times[r] = std::chrono::duration(std::chrono::high_resolution_clock::now() - t0).count(); - } - std::sort(std::begin(times), std::end(times)); - const double elapsed = times[1]; - perf_khs_out = (trials / elapsed) / 1000.0; - - crypto::hash h{}; - static_assert(sizeof(h) == crypto::HASH_SIZE, "hash size mismatch"); - std::memcpy(&h, out.data(), crypto::HASH_SIZE); - last_digest_hex_out = epee::string_tools::pod_to_hex(h); - - std::cerr << " - Hash rate: " << std::fixed << std::setprecision(2) << perf_khs_out << " kH/s\n"; -} - -// Test units -static void test_wbag_check_vectors(const std::string& path = "tests/hash/tests-wbag.txt") { - std::ifstream in(path); - if (!in) { - std::cerr << "Vectors: skipped (no " << path << ")\n"; - return; - } - - size_t cases = 0; - std::string line, pending_hash; - - auto process_pair = [&](const std::string& expect_hex, const std::string& input_hex) { - const std::string expect = trim(expect_hex); - const std::string input = trim(input_hex); - - if (expect.size() != 64) { - std::cerr << "Vectors: expected-hash len != 64\n"; std::exit(1); - } - for (char c : expect) if (!is_hex_digit(c)) { - std::cerr << "Vectors: expected-hash not hex\n"; std::exit(1); - } - - const auto msg = input.empty() ? std::vector{} : hex_to_bytes(input); - - crypto::hash out{}; - wbag_hash(msg.empty() ? nullptr : msg.data(), msg.size(), - reinterpret_cast(&out)); - - const std::string got = lower_hex(reinterpret_cast(&out), sizeof(out)); - if (!std::equal(expect.begin(), expect.end(), got.begin(), got.end(), - [](char a, char b){ - return std::tolower(static_cast(a)) == std::tolower(static_cast(b)); - })) { - std::cerr << "Vectors: mismatch at case " << (cases+1) << "\n" << " expected: " << expect << "\n" << " got: " << got << "\n"; - std::exit(1); - } - ++cases; - }; - - while (std::getline(in, line)) { - std::string s = trim(line); - if (s.empty() || s[0] == '#') continue; - - const auto sp = s.find(' '); - if (sp != std::string::npos) { - const std::string left = s.substr(0, sp); - const std::string right = s.substr(sp + 1); - process_pair(left, right); - continue; - } - - if (pending_hash.empty()) { - pending_hash = s; - continue; - } else { - process_pair(pending_hash, s); - pending_hash.clear(); - } - } - - if (!pending_hash.empty()) { - process_pair(pending_hash, ""); - pending_hash.clear(); - } - std::cerr << "Vectors: OK\n"; -} - -static void test_wbag_determinism(const cryptonote::wbag::CacheView& view) { - std::array out0{}, out1{}; - const std::vector msg(80, 0x42); - if (WBAG_hash(msg.data(), msg.size(), view.data(), out0.data()) != 0) { - std::cerr << "determinism: first hash failed\n"; std::exit(1); - } - for (int i = 0; i < 31; ++i) { - if (WBAG_hash(msg.data(), msg.size(), view.data(), out1.data()) != 0) { - std::cerr << "determinism: repeat hash failed\n"; std::exit(1); - } - if (memcmp(out0.data(), out1.data(), crypto::HASH_SIZE) != 0) { - std::cerr << "determinism: mismatch on repeat " << i << "\n"; std::exit(1); - } - } - std::cerr << " - Determinism: OK\n"; -} - -static void test_wbag_edge_inputs(const cryptonote::wbag::CacheView& view) { - std::array out{}; - if (WBAG_hash(nullptr, 0, view.data(), out.data()) != 0) { - std::cerr << "edge: len=0 failed\n"; std::exit(1); - } - uint8_t buf[66]; for (int i = 0; i < 66; ++i) buf[i] = static_cast(i); - const size_t lens[] = {1, 63, 64, 65}; - for (size_t L : lens) { - if (WBAG_hash(buf + 1, L, view.data(), out.data()) != 0) { - std::cerr << "edge: len=" << L << " failed\n"; std::exit(1); - } - } - std::cerr << " - Edge inputs: OK\n"; -} - -static void test_wbag_api_errors(const cryptonote::wbag::CacheView& view) { - std::array out{}; - errno = 0; - if (WBAG_hash(nullptr, 4, view.data(), out.data()) != -1 || errno != EINVAL) { - std::cerr << "api: expected EINVAL for bd=null,len>0\n"; std::exit(1); - } - errno = 0; - if (WBAG_hash(nullptr, 0, nullptr, out.data()) != -1 || errno != EINVAL) { - std::cerr << "api: expected EINVAL for cache=null\n"; std::exit(1); - } - errno = 0; - if (WBAG_hash(nullptr, 0, view.data(), nullptr) != -1 || errno != EINVAL) { - std::cerr << "api: expected EINVAL for out32=null\n"; std::exit(1); - } - std::cerr << " - API errors check: OK\n"; -} - -static void test_wbag_cache_sha_stability(uint64_t height, - const std::string& sha_hex_again) { - static std::unordered_map seen; - auto it = seen.find(height); - if (it == seen.end()) { - seen.emplace(height, sha_hex_again); - } else if (it->second != sha_hex_again) { - std::cerr << "cache-sha: instability for height " << height << "\n"; std::exit(1); - } - for (const auto& kv : seen) { - if (kv.first != height && kv.second == sha_hex_again) { - std::cerr << "cache-sha: SHA collision across anchors " << kv.first << " and " << height << "\n"; std::exit(1); - } - } - std::cerr << " - Cache stability/distinctness: OK\n"; -} - -static void test_wbag_cache_hash_check(uint64_t height, const std::string& sha256_hex) { - auto it = kGoldenCacheSHA256.find(height); - if (it == kGoldenCacheSHA256.end()) { - std::cerr << " - Cache SHA-256 hash match: skipped (no golden)\n"; - return; - } - if (sha256_hex != it->second) { - std::cerr << " - Cache SHA-256 hash match: MISMATCH\n" - << " expected: " << it->second << "\n" - << " got: " << sha256_hex << "\n"; - std::exit(1); - } - std::cerr << " - Cache SHA-256 hash match: OK\n"; -} - -static void test_wbag_fuzz(const cryptonote::wbag::CacheView& view, uint32_t iters=2000) { - std::vector buf(1u<<20); - std::mt19937_64 rng(0xC0FFEEULL); - std::uniform_int_distribution len_dist(0, buf.size()); - std::uniform_int_distribution off_dist(0, 63); - - for (uint32_t i=0; i(rng()); - size_t L = len_dist(rng); - size_t off = std::min(off_dist(rng), L); - std::array out{}; - if (WBAG_hash(buf.data()+off, L-off, view.data(), out.data()) != 0) { - std::cerr << "fuzz: hash failed at iter " << i << "\n"; std::exit(1); - } - } - std::cerr << " - Fuzz: OK\n"; -} - -static void test_wbag_avalanche(const cryptonote::wbag::CacheView& view) { - std::array msg{}; - std::array base{}, alt{}; - if (WBAG_hash(msg.data(), msg.size(), view.data(), base.data()) != 0) { - std::cerr << "avalanche: base hash failed\n"; std::exit(1); - } - auto hdist = [](const uint8_t* a, const uint8_t* b){ - unsigned c=0; - for (size_t i=0;i>3] ^= (1u<<(bit&7)); - if (WBAG_hash(msg.data(), msg.size(), view.data(), alt.data()) != 0) { - std::cerr << "avalanche: alt hash failed at bit " << bit << "\n"; std::exit(1); - } - sum += hdist(base.data(), alt.data()); - msg[bit>>3] ^= (1u<<(bit&7)); - } - double avg = sum / (64*8); - if (avg < 112 || avg > 144) { - std::cerr << "avalanche: weak diffusion, avg bit flip=" << avg << "\n"; std::exit(1); - } - std::cerr << " - Avalanche: OK (avg " << std::fixed << std::setprecision(1) << avg << " bits)\n"; -} - -static void test_wbag_unaligned_inputs(const cryptonote::wbag::CacheView& view) { - std::vector buf(4096); - for (size_t i = 0; i < buf.size(); ++i) buf[i] = uint8_t(i); - const size_t L = 1024, off = 37; // misaligned - std::array a{}, b{}; - - if (WBAG_hash(buf.data() + off, L, view.data(), a.data()) != 0) { - std::cerr << "unaligned: hash(slice) failed\n"; std::exit(1); - } - std::vector tmp(buf.begin() + off, buf.begin() + off + L); - if (WBAG_hash(tmp.data(), tmp.size(), view.data(), b.data()) != 0) { - std::cerr << "unaligned: hash(copy) failed\n"; std::exit(1); - } - if (std::memcmp(a.data(), b.data(), crypto::HASH_SIZE) != 0) { - std::cerr << "unaligned: mismatch\n"; std::exit(1); - } - std::cerr << " - Unaligned inputs: OK\n"; -} - -static void test_wbag_large_sizes_smoke(const cryptonote::wbag::CacheView& view) { - const size_t sizes[] = {0,1,63,64,65,1024, 65536, 1048576, 1048583}; - std::vector buf(1048583); - for (size_t i = 0; i < buf.size(); ++i) buf[i] = uint8_t((i * 1315423911u) >> 24); - - std::array out{}, out2{}; - for (size_t s : sizes) { - if (WBAG_hash(buf.data(), s, view.data(), out.data()) != 0) { - std::cerr << "large: hash failed at size " << s << "\n"; std::exit(1); - } - } - // Diffusion smoke: flip byte and expect output change - const size_t pos = 1048576 - 1; - if (WBAG_hash(buf.data(), 1048576, view.data(), out.data()) != 0) { - std::cerr << "large: baseline hash failed\n"; std::exit(1); - } - buf[pos] ^= 0x80; - if (WBAG_hash(buf.data(), 1048576, view.data(), out2.data()) != 0) { - std::cerr << "large: alt hash failed\n"; std::exit(1); - } - if (std::memcmp(out.data(), out2.data(), crypto::HASH_SIZE) == 0) { - std::cerr << "large: no change after byte flip\n"; std::exit(1); - } - std::cerr << " - Large-size smoke: OK\n"; -} - -static void test_wbag_multithread_smoke(const cryptonote::wbag::CacheView& view) { - const int T = 8, I = 256; - std::vector common(80, 0x42); - std::array ref{}; - if (WBAG_hash(common.data(), common.size(), view.data(), ref.data()) != 0) { - std::cerr << "threaded: ref hash failed\n"; std::exit(1); - } - - std::atomic err{0}; +// ---------------------- CPU bench ---------------------- +static double bench_cpu_once(const WbagCacheView& view, int threads, int trials) { std::vector th; - th.reserve(T); - for (int t = 0; t < T; ++t) { - th.emplace_back([&, t]{ - std::vector msg(1000 + t); - for (size_t i = 0; i < msg.size(); ++i) msg[i] = uint8_t((i + t) & 0xFF); - std::array a{}, b{}, c{}; - for (int i = 0; i < I; ++i) { - if (WBAG_hash(msg.data(), msg.size(), view.data(), a.data()) != 0) { ++err; return; } - if (WBAG_hash(msg.data(), msg.size(), view.data(), b.data()) != 0) { ++err; return; } - if (std::memcmp(a.data(), b.data(), crypto::HASH_SIZE) != 0) { ++err; return; } - if (WBAG_hash(common.data(), common.size(), view.data(), c.data()) != 0) { ++err; return; } - if (std::memcmp(ref.data(), c.data(), crypto::HASH_SIZE) != 0) { ++err; return; } - } - }); - } - for (auto& x : th) x.join(); - if (err.load() != 0) { std::cerr << "threaded: mismatch\n"; std::exit(1); } - std::cerr << " - Multithread smoke: OK\n"; -} + th.reserve(std::max(1, threads)); -static void test_wbag_first_use_race() { - constexpr int T = 16; + std::atomic ready{0}; std::atomic go{false}; - std::atomic err{0}; - std::vector th; th.reserve(T); - for (int i=0;i out{}; - uint8_t b = 0x42; - try { wbag_hash(&b, 1, reinterpret_cast(out.data())); } - catch (...) { ++err; } - }); - } - go.store(true, std::memory_order_release); - for (auto& x: th) x.join(); - if (err.load() != 0) { std::cerr << "first-use race: fail\n"; std::exit(1); } - std::cerr << "\nFirst use: OK\n"; -} + const std::vector base_msg(80, 0x42); -static void test_wbag_geometry() { - const uint64_t cb = WBAG_cache_bytes(); - const uint32_t eb = WBAG_entry_bytes(); - if (cb == 0 || eb == 0 || (cb % eb) != 0) { - std::cerr << "geometry: invalid\n"; std::exit(1); - } - std::cerr << "Geometry: OK\n"; -} - -static void run_wbag_tests(const cryptonote::wbag::CacheView& view, uint64_t height, const std::string& sha256_hex, const std::vector& block) { - if (!sha256_hex.empty()) { - test_wbag_cache_sha_stability(height, sha256_hex); - test_wbag_cache_hash_check(height, sha256_hex); - } - test_wbag_determinism(view); - test_wbag_edge_inputs(view); - test_wbag_api_errors(view); - test_wbag_fuzz(view); - test_wbag_avalanche(view); - test_wbag_unaligned_inputs(view); - test_wbag_large_sizes_smoke(view); - test_wbag_multithread_smoke(view); -} - -// Start WBG tests -int wbag_test() { - struct Row { - uint64_t height = 0; - uint64_t epoch = 0; - size_t cache_size = 0; - double ensure_sec = 0.0; - double perf_khs = 0.0; - std::string sha256_hex; - std::string last_digest_hex; - std::string note; + auto worker = [&](int tid) { + std::vector msg = base_msg; + if (!msg.empty()) msg[0] ^= uint8_t(tid * 31); + std::array out{}; + ready.fetch_add(1, std::memory_order_acq_rel); + while (!go.load(std::memory_order_acquire)) {} + for (int i = 0; i < trials; ++i) { + WBAG_hash(msg.data(), msg.size(), + view.cache, + view.resident, view.resident_size, + out.data()); + } }; - // Params - const int trials = 2048; - const std::vector block(80, 0x42); - const size_t cache_bytes = WBAG_cache_bytes(); - const double cache_mb = static_cast(cache_bytes) / (1024.0 * 1024.0); - const size_t sp_bytes = WBAG_scratchpad(); - const double sp_mb = static_cast(sp_bytes) / (1024.0 * 1024.0); - const int total = static_cast(sizeof(k_wbag_anchor_samples) / sizeof(k_wbag_anchor_samples[0])); - uint32_t sig = WBAG_params_signature(); + for (int i = 0; i < threads; ++i) th.emplace_back(worker, i); + while (ready.load(std::memory_order_acquire) != threads) {} + const auto t1 = std::chrono::high_resolution_clock::now(); + go.store(true, std::memory_order_release); + for (auto& x : th) x.join(); - std::cerr << "Params:" - << " sig=0x" << std::hex << sig << std::dec - << "\nWBAG_CACHE_BYTES=" << cache_mb - << ", WBAG_ENTRY_BYTES=" << WBAG_entry_bytes() - << ", WBAG_SCRATCHPAD=" << sp_mb << "," - << "\nWBAG_ROUNDS=" << WBAG_rounds() - << ", WBAG_ARGON2_P=" << WBAG_argon_p() - << ", WBAG_ARGON2_T=" << WBAG_argon_t(); - - std::cerr << "\n\nStarting tests...\n"; - - test_wbag_first_use_race(); - test_wbag_geometry(); - test_wbag_check_vectors(); - - std::cerr << "Hashing " << total << " rounds...\n"; - - // Report - std::ostringstream report; - report << "\n\n"; - report << std::left << std::setw(12) << "Height" << std::setw(66) << "SHA256 (cache)" << "Digest\n"; - report << std::string(12+66+64, '-') << "\n"; - - for (int idx = 0; idx < total; ++idx) { - const auto &sample = k_wbag_anchor_samples[idx]; - Row r; - r.height = sample.height; - r.epoch = cryptonote::wbag::epoch(sample.height); - crypto::hash anchor{}; - if (!epee::string_tools::hex_to_pod(sample.block_id_hex, anchor)) { - r.note = "bad anchor hex"; - } else { - std::cerr << "[" << (idx+1) << "/" << total << "] height " << r.height << " epoch " << r.epoch << ": generating cache..." << std::flush; - - if (idx == 0 && total > 1) { - crypto::hash other{}; - if (epee::string_tools::hex_to_pod(k_wbag_anchor_samples[1].block_id_hex, other)) { - const uint64_t other_ep = cryptonote::wbag::epoch(k_wbag_anchor_samples[1].height); - (void)cryptonote::wbag::ensure_cache(other_ep, other); - } - } - - const auto t0 = std::chrono::high_resolution_clock::now(); - if (!cryptonote::wbag::ensure_cache(r.epoch, anchor)) { - r.note = "ensure_cache failed"; - std::cerr << " FAILED\n"; - } else { - r.ensure_sec = std::chrono::duration(std::chrono::high_resolution_clock::now() - t0).count(); - std::cerr << " " << std::fixed << std::setprecision(3) << r.ensure_sec << " s\n"; - auto view = cryptonote::wbag::cache_view(); - if (!view || !view.data() || view.size() == 0) { - r.note = "cache_view invalid"; - } else { - r.cache_size = view.size(); - - if (!write_cache_and_sha256(view, r.sha256_hex, r.note)) { - } - - // Run tests - run_wbag_tests(view, r.height, r.sha256_hex, block); - - // Benchmark - if (r.note.empty()) { - benchmark_hash(view, trials, block, r.perf_khs, r.last_digest_hex); - } - } - } - } - report << std::left - << std::setw(12) << (r.height ? std::to_string(r.height) : std::string("-")) - << std::setw(66) << (r.sha256_hex.empty() ? std::string("-") : r.sha256_hex) - << (r.last_digest_hex.empty() ? std::string("-") : r.last_digest_hex) - << (r.note.empty() ? "" : (" [" + r.note + "]")) - << "\n"; - } - std::cout << report.str(); - return 0; + const double sec = std::max(1e-12, + std::chrono::duration(std::chrono::high_resolution_clock::now() - t1).count()); + return (double(trials) * double(threads)) / sec; } +// ---------------------- Public C API wrappers ---------------------- +extern "C" void wbag_hash(const void* data, size_t length, char* hash) { + static std::once_flag once; + static bool ok = false; + static cryptonote::wbag::CacheView g_cache; + static std::vector g_resident; + static size_t g_resident_size = 0; + + std::call_once(once, [&]{ + const auto& a = k_wbag_anchor_samples[0]; + crypto::hash anchor{}; + if (!epee::string_tools::hex_to_pod(a.block_id_hex, anchor)) return; + const uint64_t ep = cryptonote::wbag::epoch(a.height); + if (!cryptonote::wbag::ensure_cache(ep, anchor)) return; + g_cache = cryptonote::wbag::cache_view(); + if (!g_cache || !g_cache.data() || g_cache.size() == 0) return; + + g_resident_size = WBAG_scratchpad_resident(); + g_resident.resize(g_resident_size); + + const uint64_t rseed = WBAG_resident_seed_from_cache(g_cache.data(), g_cache.size()); + if (WBAG_resident_fill(g_resident.data(), g_resident.size(), rseed) != 0) return; + + ok = true; + }); + + uint8_t out[32]{}; + if (!ok) { std::memset(hash, 0, 32); return; } + + if (WBAG_hash(data, length, g_cache.data(), g_resident.data(), g_resident_size, out) != 0) { + std::memset(hash, 0, 32); + return; + } + std::memcpy(hash, out, 32); +} + +// ---------------------- CLI ---------------------- extern "C" int wbag_hashrate_cli(int argc, char** argv) { - // Args bool want_cpu = false, want_gpu = false; int threads = std::max(1u, std::thread::hardware_concurrency()); - int runs = 5; - int trials = 2048; - bool regen = false; + int trials = 2048; + int runs = 1; + int anchor_index = 0; for (int i = 1; i < argc; ++i) { const std::string a = argv[i]; @@ -633,101 +165,47 @@ extern "C" int wbag_hashrate_cli(int argc, char** argv) { else if (a == "--cpu") want_cpu = true; else if (a == "--gpu") want_gpu = true; else if (a == "--threads" && i + 1 < argc) { threads = std::max(1, std::atoi(argv[++i])); } - else if (a == "--runs" && i + 1 < argc) { runs = std::max(1, std::atoi(argv[++i])); } - else if (a == "--trials" && i + 1 < argc) { trials = std::max(1, std::atoi(argv[++i])); } - else if (a == "--regen") { regen = true; } + else if (a == "--trials" && i + 1 < argc) { trials = std::max(1, std::atoi(argv[++i])); } + else if (a == "--runs" && i + 1 < argc) { runs = std::max(1, std::atoi(argv[++i])); } + else if (a == "--anchor" && i + 1 < argc) { anchor_index = std::atoi(argv[++i]); } } if (!(want_cpu || want_gpu)) want_cpu = true; + if (anchor_index < 0) anchor_index = 0; + if (anchor_index >= k_wbag_anchor_count) anchor_index = k_wbag_anchor_count - 1; - // Pick an anchor to force cache build; flip if --regen to avoid reusing an already-built cache - const int anchor_idx = regen ? 1 : 0; - if (anchor_idx >= int(sizeof(k_wbag_anchor_samples)/sizeof(k_wbag_anchor_samples[0]))) { - std::cerr << "No anchors available\n"; - return 1; - } - const auto &sample = k_wbag_anchor_samples[anchor_idx]; - - crypto::hash anchor{}; - if (!epee::string_tools::hex_to_pod(sample.block_id_hex, anchor)) { - std::cerr << "Bad anchor hex\n"; - return 1; - } - const uint64_t ep = cryptonote::wbag::epoch(sample.height); - - // Cache gen timing - auto t0 = std::chrono::high_resolution_clock::now(); - if (!cryptonote::wbag::ensure_cache(ep, anchor)) { + WbagCacheView cpu_view{}; + cryptonote::wbag::CacheView gpu_view{}; + double cache_sec = 0.0; + if (!ensure_wbag_cache(anchor_index, cpu_view, gpu_view, cache_sec)) { std::cerr << "Cache build failed\n"; return 1; } - const double cache_sec = - std::chrono::duration(std::chrono::high_resolution_clock::now() - t0).count(); - auto view = cryptonote::wbag::cache_view(); - if (!view || !view.data() || view.size() == 0) { - std::cerr << "Cache view invalid\n"; - return 1; - } - - auto print_rate = [](double hs) { - std::ostringstream o; - if (hs >= 1000.0) { - o << std::fixed << std::setprecision(2) << (hs / 1000.0) << " kH/s"; - } else { - o << std::fixed << std::setprecision(0) << hs << " H/s"; - } - return o.str(); - }; - - // Header - std::cout << "Miner: " << (want_cpu ? "CPU" : "GPU") << "\n"; - std::cout << "Cache Gen Time: " << std::fixed << std::setprecision(2) << cache_sec << " s\n"; - std::cout << "Threads: " << threads << "\n\n"; - std::cout << "Hashing...\n\n"; - - // Common msg - std::vector base_msg(80, 0x42); + std::cout << "Mode: " << (want_cpu ? "CPU" : "GPU") << "\n"; + std::cout << "Anchor: [" << anchor_index << "] height=" << k_wbag_anchor_samples[anchor_index].height << "\n"; + std::cout << "Cache Gen Time: " << std::fixed << std::setprecision(3) << cache_sec << " s\n"; + std::cout << "Threads: " << threads << "\n"; + std::cout << "Trials: " << trials << "\n\n"; + std::cout << "Hashing...\n"; if (want_cpu) { for (int r = 1; r <= runs; ++r) { - std::atomic ready{0}; - std::atomic go{false}; - std::vector th; th.reserve(threads); - - auto worker = [&](int tid) { - // Slightly vary input per-thread to avoid unreal cache friendliness - std::vector msg = base_msg; - if (!msg.empty()) msg[0] ^= uint8_t(tid * 31); - ready.fetch_add(1, std::memory_order_acq_rel); - while (!go.load(std::memory_order_acquire)) {} - std::array out{}; - for (int i = 0; i < trials; ++i) { - WBAG_hash(msg.data(), msg.size(), view.data(), out.data()); - } - }; - - for (int i = 0; i < threads; ++i) th.emplace_back(worker, i); - while (ready.load(std::memory_order_acquire) != threads) {} - auto t1 = std::chrono::high_resolution_clock::now(); - go.store(true, std::memory_order_release); - for (auto &x : th) x.join(); - const double sec = std::chrono::duration(std::chrono::high_resolution_clock::now() - t1).count(); - const double hs = (double(trials) * double(threads)) / std::max(sec, 1e-12); + const double hs = bench_cpu_once(cpu_view, threads, trials); std::cout << "[" << r << "/" << runs << "] HR: " << print_rate(hs) << "\n"; } - std::cout << "\n"; return 0; } - // GPU path #if defined(WBAG_HAVE_WBAG_GPU_BENCH) for (int r = 1; r <= runs; ++r) { double hs = 0.0; - const int rc = wbag_gpu_bench_once(view, threads, trials, hs); - if (rc != 0) { std::cerr << "GPU bench failed (rc=" << rc << ")\n"; return 1; } + const int rc = wbag_gpu_bench_once(gpu_view, threads, trials, hs); + if (rc != 0) { + std::cerr << "GPU bench failed (rc=" << rc << ")\n"; + return 1; + } std::cout << "[" << r << "/" << runs << "] HR: " << print_rate(hs) << "\n"; } - std::cout << "\n"; return 0; #else std::cerr << "GPU support not compiled. Define WBAG_HAVE_WBAG_GPU_BENCH and provide wbag_gpu_bench_once().\n"; diff --git a/tests/hash/wbag_test.h b/tests/hash/wbag_test.h index 3905bb913..ef8ebe1c0 100644 --- a/tests/hash/wbag_test.h +++ b/tests/hash/wbag_test.h @@ -1,6 +1,40 @@ #pragma once +#include #include +struct WbagAnchorSample { + uint64_t height; + const char* block_id_hex; // 64 hex chars +}; + +extern const WbagAnchorSample k_wbag_anchor_samples[]; +extern const int k_wbag_anchor_count; + +// CPU bench view (adds resident buffer) +struct WbagCacheView { + const uint8_t* cache = nullptr; + size_t cache_size = 0; + const uint8_t* resident = nullptr; + size_t resident_size = 0; +}; + +#ifdef __cplusplus +extern "C" { +#endif + +// Expose wrapper so main.cpp can reference {"wbag", wbag_hash} void wbag_hash(const void* data, size_t length, char* hash); -int wbag_test(); +// CLI entry used by main.cpp via --wbag-hr +int wbag_hashrate_cli(int argc, char** argv); + +#ifdef __cplusplus +} // extern "C" +#endif + +#if defined(WBAG_HAVE_WBAG_GPU_BENCH) +namespace cryptonote { namespace wbag { struct CacheView; } } +extern "C" int wbag_gpu_bench_once(const cryptonote::wbag::CacheView& view, + int threads, int trials, + double& hashes_per_sec_out); +#endif