codeberg-wowario/wownero
1
0
Fork 0
mirror of https://codeberg.org/wowario/wownero synced 2026-10-17 05:41:40 +02:00
Code Issues Packages Projects Releases Wiki Activity

fix

Browse Source
This commit is contained in:
wowario
2025-09-26 00:22:05 +03:00
parent 0f5b395f83
commit e5119dd2ea
7 changed files with 484 additions and 841 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/crypto/wbag-hash.c
+116 -60
View File
@@ -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"
src/crypto/wbag-hash.h
+7 -1
View File
@@ -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);
tests/hash/CMakeLists.txt
-3
View File
@@ -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")
tests/hash/main.cpp
-3
View File
@@ -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);
tests/hash/wbag_gpu_bench.cu
+193 -118
View File
@@ -1,197 +1,272 @@
#include <cuda_runtime.h>
#include <cstdint>
#include <vector>
#include <algorithm>
#include <chrono>
#include <cstdio>
#include <cstring>
#include <algorithm>
#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<<<grid, block>>>(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<<<grid, block>>>(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<<<grid, block>>>(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<uint64_t> h_out;
h_out.resize(static_cast<size_t>(threads));
if (auto e = cudaMemcpy(h_out.data(), d_out,
sizeof(uint64_t) * static_cast<size_t>(threads),
cudaMemcpyDeviceToHost);
e != cudaSuccess) {
fprintf(stderr, "GPU memcpy failed: %s\n", cudaGetErrorString(e));
cudaFree(d_out); cudaFree(d_scratch);
// Sanity
std::vector<uint64_t> 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;
}
tests/hash/wbag_test.cpp
+133 -655
View File
@@ -1,631 +1,163 @@
#include <array>
#include <cerrno>
#include <atomic>
#include <chrono>
#include <cstddef>
#include <cstdio>
#include <cstdint>
#include <cstdlib>
#include <cstring>
#include <fstream>
#include <iomanip>
#include <ios>
#include <iostream>
#include <mutex>
#include <sstream>
#include <string>
#include <unordered_map>
#include <thread>
#include <vector>
#include <algorithm>
#include <cctype>
#include <stdexcept>
#include <thread>
#include <atomic>
#include <random>
#include <mutex>
#if defined(_MSC_VER) && !defined(__clang__)
#include <intrin.h>
#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<uint64_t, std::string> 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<uint8_t> 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<uint8_t> 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<uint8_t>((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<double>(std::chrono::high_resolution_clock::now() - t0).count();
if (WBAG_hash(data, length, cached.data(), reinterpret_cast<uint8_t*>(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<const char*>(view.data()), static_cast<std::streamsize>(view.size()));
if (!out.good()) { note = "write cache bytes failed"; return false; }
}
const size_t resident_bytes = WBAG_scratchpad_resident();
static std::vector<uint8_t> 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<uint8_t>& block, double& perf_khs_out, std::string& last_digest_hex_out) {
std::vector<uint8_t> 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<double>(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<uint8_t>{} : hex_to_bytes(input);
crypto::hash out{};
wbag_hash(msg.empty() ? nullptr : msg.data(), msg.size(),
reinterpret_cast<char*>(&out));
const std::string got = lower_hex(reinterpret_cast<const uint8_t*>(&out), sizeof(out));
if (!std::equal(expect.begin(), expect.end(), got.begin(), got.end(),
[](char a, char b){
return std::tolower(static_cast<unsigned char>(a)) == std::tolower(static_cast<unsigned char>(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<uint8_t, crypto::HASH_SIZE> out0{}, out1{};
const std::vector<uint8_t> 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<uint8_t, crypto::HASH_SIZE> 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<uint8_t>(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<uint8_t, crypto::HASH_SIZE> 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<uint64_t, std::string> 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<uint8_t> buf(1u<<20);
std::mt19937_64 rng(0xC0FFEEULL);
std::uniform_int_distribution<size_t> len_dist(0, buf.size());
std::uniform_int_distribution<size_t> off_dist(0, 63);
for (uint32_t i=0; i<iters; ++i) {
for (auto& b : buf) b = static_cast<uint8_t>(rng());
size_t L = len_dist(rng);
size_t off = std::min(off_dist(rng), L);
std::array<uint8_t, crypto::HASH_SIZE> 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<uint8_t,64> msg{};
std::array<uint8_t, crypto::HASH_SIZE> 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<crypto::HASH_SIZE;i++) c += pc8(unsigned(a[i]^b[i]));
return c;
};
double sum=0;
for (int bit=0; bit<64*8; ++bit) {
msg[bit>>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<uint8_t> 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<uint8_t, crypto::HASH_SIZE> 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<uint8_t> 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<uint8_t> buf(1048583);
for (size_t i = 0; i < buf.size(); ++i) buf[i] = uint8_t((i * 1315423911u) >> 24);
std::array<uint8_t, crypto::HASH_SIZE> 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<uint8_t> common(80, 0x42);
std::array<uint8_t, crypto::HASH_SIZE> 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<int> err{0};
// ---------------------- CPU bench ----------------------
static double bench_cpu_once(const WbagCacheView& view, int threads, int trials) {
std::vector<std::thread> th;
th.reserve(T);
for (int t = 0; t < T; ++t) {
th.emplace_back([&, t]{
std::vector<uint8_t> msg(1000 + t);
for (size_t i = 0; i < msg.size(); ++i) msg[i] = uint8_t((i + t) & 0xFF);
std::array<uint8_t, crypto::HASH_SIZE> 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<int> ready{0};
std::atomic<bool> go{false};
std::atomic<int> err{0};
std::vector<std::thread> th; th.reserve(T);
for (int i=0;i<T;++i) {
th.emplace_back([&]{
while (!go.load(std::memory_order_acquire)) {}
std::array<uint8_t, crypto::HASH_SIZE> out{};
uint8_t b = 0x42;
try { wbag_hash(&b, 1, reinterpret_cast<char*>(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<uint8_t> 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<uint8_t>& 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<uint8_t> msg = base_msg;
if (!msg.empty()) msg[0] ^= uint8_t(tid * 31);
std::array<uint8_t, 32> 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<uint8_t> block(80, 0x42);
const size_t cache_bytes = WBAG_cache_bytes();
const double cache_mb = static_cast<double>(cache_bytes) / (1024.0 * 1024.0);
const size_t sp_bytes = WBAG_scratchpad();
const double sp_mb = static_cast<double>(sp_bytes) / (1024.0 * 1024.0);
const int total = static_cast<int>(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<double>(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<double>(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<uint8_t> 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<double>(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<uint8_t> 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<int> ready{0};
std::atomic<bool> go{false};
std::vector<std::thread> th; th.reserve(threads);
auto worker = [&](int tid) {
// Slightly vary input per-thread to avoid unreal cache friendliness
std::vector<uint8_t> 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<uint8_t, crypto::HASH_SIZE> 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<double>(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";
tests/hash/wbag_test.h
+35 -1
View File
@@ -1,6 +1,40 @@
#pragma once
#include <cstdint>
#include <cstddef>
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