283 lines
17 KiB
Text
283 lines
17 KiB
Text
// Attack-pass F9: the header-grinding rate measurement on one NVIDIA card. NOT a miner: no pool, no network, no wallet.
|
|
//
|
|
// Built against a memory-hard class v4 pack (program.h, vectors.h, kernel.cu, kernel_bound.cu) plus f9-variants.cu
|
|
// (make-variants.py). Fills the cache and the dataset with the pack's own kernels, checks both against vectors.h,
|
|
// checks the bound hash against igneum-pow's reference lines (--ref), checks the per-warp kernel reproduces the honest
|
|
// kernel on an all-equal table (known-pass) and that the forced kernels do not (known-fail), then times the variants
|
|
// in interleaved rounds. Every phase prints its wall-clock window in epoch milliseconds so the nvidia-smi power log can
|
|
// be attributed to it (join-power.py).
|
|
//
|
|
// f9-host --prehash <64 hex> --ref ref.txt --table-random f --table-k10 f --table-k14 f [--rounds 5] [--seconds 8]
|
|
// [--batch-log2 24] [--block-warps 1]
|
|
#include <cuda_runtime.h>
|
|
#include <cstdint>
|
|
#include <cstdio>
|
|
#include <cstdlib>
|
|
#include <cstring>
|
|
#include <chrono>
|
|
#include <cmath>
|
|
#include <fstream>
|
|
#include <string>
|
|
#include <vector>
|
|
#include "program.h"
|
|
#include "vectors.h"
|
|
|
|
#define CUDA_CHECK(call) do { cudaError_t err_ = (call); if (err_ != cudaSuccess) { \
|
|
std::fprintf(stderr, "CUDA error: %s (%d) at %s:%d in %s\n", cudaGetErrorString(err_), (int)err_, __FILE__, __LINE__, #call); \
|
|
std::exit(2); } } while (0)
|
|
|
|
struct IgneumInitWords { uint32_t w[8]; };
|
|
cudaError_t igneum_launch_hash_bound(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask, IgneumInitWords iw, uint32_t nonces, uint32_t blockWarps);
|
|
cudaError_t f9_launch_perwarp(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask, const IgneumInitWords* tbl, uint32_t nonces, uint32_t blockWarps);
|
|
cudaError_t f9_launch_forced4(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask, IgneumInitWords iw, uint32_t nonces, uint32_t blockWarps);
|
|
cudaError_t f9_launch_forcedall(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask, IgneumInitWords iw, uint32_t nonces, uint32_t blockWarps);
|
|
cudaError_t f9_launch_forced1(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask, IgneumInitWords iw, uint32_t nonces, uint32_t blockWarps);
|
|
cudaError_t f9_launch_pair(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask, IgneumInitWords iw, uint32_t nonces, uint32_t blockWarps);
|
|
|
|
static uint64_t fnv1a64(const void* p, size_t n) {
|
|
const uint8_t* b = (const uint8_t*)p;
|
|
uint64_t h = 0xcbf29ce484222325ull;
|
|
for (size_t i = 0; i < n; ++i) { h ^= b[i]; h *= 0x100000001b3ull; }
|
|
return h;
|
|
}
|
|
static void seedWordsFromBytes(const uint8_t* b, size_t n, uint32_t out[8]) {
|
|
for (uint64_t salt = 0; salt < 4; ++salt) {
|
|
uint64_t h = 0xcbf29ce484222325ull ^ (salt * 0x9E3779B97F4A7C15ull);
|
|
for (size_t i = 0; i < n; ++i) { h ^= b[i]; h *= 0x100000001b3ull; }
|
|
h ^= h >> 33; h *= 0xff51afd7ed558ccdull; h ^= h >> 33;
|
|
out[2 * salt] = (uint32_t)h;
|
|
out[2 * salt + 1] = (uint32_t)(h >> 32);
|
|
}
|
|
}
|
|
static IgneumInitWords blockInit(const uint8_t prehash[32], uint32_t nonceHi) {
|
|
uint8_t b[49];
|
|
std::memcpy(b, "igneum-block/", 13);
|
|
std::memcpy(b + 13, prehash, 32);
|
|
b[45] = (uint8_t)nonceHi; b[46] = (uint8_t)(nonceHi >> 8); b[47] = (uint8_t)(nonceHi >> 16); b[48] = (uint8_t)(nonceHi >> 24);
|
|
IgneumInitWords iw;
|
|
seedWordsFromBytes(b, 49, iw.w);
|
|
return iw;
|
|
}
|
|
static bool unhex(const std::string& s, std::vector<uint8_t>& out) {
|
|
if (s.size() % 2) return false;
|
|
out.clear();
|
|
for (size_t i = 0; i < s.size(); i += 2) out.push_back((uint8_t)std::strtoul(s.substr(i, 2).c_str(), nullptr, 16));
|
|
return true;
|
|
}
|
|
static double nowMs() { return std::chrono::duration<double, std::milli>(std::chrono::steady_clock::now().time_since_epoch()).count(); }
|
|
static long long epochMs() { return std::chrono::duration_cast<std::chrono::milliseconds>(std::chrono::system_clock::now().time_since_epoch()).count(); }
|
|
|
|
static std::vector<IgneumInitWords> loadTable(const std::string& path, size_t warps) {
|
|
std::vector<IgneumInitWords> t(warps);
|
|
std::ifstream f(path, std::ios::binary);
|
|
if (!f) { std::fprintf(stderr, "cannot open table %s\n", path.c_str()); std::exit(2); }
|
|
f.read((char*)t.data(), (std::streamsize)(warps * sizeof(IgneumInitWords)));
|
|
if ((size_t)f.gcount() != warps * sizeof(IgneumInitWords)) { std::fprintf(stderr, "table %s short: %lld bytes\n", path.c_str(), (long long)f.gcount()); std::exit(2); }
|
|
return t;
|
|
}
|
|
|
|
enum Variant { HONEST, PERWARP_RANDOM, PERWARP_K10, PERWARP_K14, PAIR, FORCED1, FORCED4, FORCEDALL, NVARIANTS };
|
|
static const char* NAMES[NVARIANTS] = { "honest", "perwarp-random", "perwarp-k10", "perwarp-k14", "pair", "forced1", "forced4", "forcedall" };
|
|
|
|
int main(int argc, char** argv) {
|
|
std::string prehashHex(64, '0'), refPath, tRandom, tK10, tK14;
|
|
int rounds = 5, seconds = 8, batchLog2 = 24, blockWarps = 1;
|
|
for (int i = 1; i < argc; ++i) {
|
|
std::string a = argv[i];
|
|
auto next = [&]() -> std::string { if (i + 1 >= argc) { std::fprintf(stderr, "missing value for %s\n", a.c_str()); std::exit(2); } return argv[++i]; };
|
|
if (a == "--prehash") prehashHex = next();
|
|
else if (a == "--ref") refPath = next();
|
|
else if (a == "--table-random") tRandom = next();
|
|
else if (a == "--table-k10") tK10 = next();
|
|
else if (a == "--table-k14") tK14 = next();
|
|
else if (a == "--rounds") rounds = std::atoi(next().c_str());
|
|
else if (a == "--seconds") seconds = std::atoi(next().c_str());
|
|
else if (a == "--batch-log2") batchLog2 = std::atoi(next().c_str());
|
|
else if (a == "--block-warps") blockWarps = std::atoi(next().c_str());
|
|
else { std::fprintf(stderr, "unknown argument %s\n", a.c_str()); return 2; }
|
|
}
|
|
std::vector<uint8_t> ph;
|
|
if (!unhex(prehashHex, ph) || ph.size() != 32) { std::fprintf(stderr, "--prehash needs 64 hex\n"); return 2; }
|
|
cudaDeviceProp prop; std::memset(&prop, 0, sizeof prop);
|
|
CUDA_CHECK(cudaGetDeviceProperties(&prop, 0));
|
|
std::printf("f9-host: %s, %d SMs, %.0f MiB, pack %s class %s\n", prop.name, prop.multiProcessorCount, prop.totalGlobalMem / 1048576.0, IGNEUM_SEED_STRING, IGNEUM_PROGRAM_CLASS);
|
|
|
|
// 1. cache: filled on the GPU by the pack's kernel, checked against the pack's FNV, head and last line
|
|
const uint32_t cacheWords = 1u << IGNEUM_CACHE_LOG2_WORDS;
|
|
uint32_t* dCache = nullptr;
|
|
CUDA_CHECK(cudaMalloc((void**)&dCache, (size_t)cacheWords * 4));
|
|
double t0 = nowMs();
|
|
CUDA_CHECK(igneum_launch_cache_fill(dCache, IGNEUM_CACHE_SEGMENTS));
|
|
CUDA_CHECK(cudaDeviceSynchronize());
|
|
double fillMs = nowMs() - t0;
|
|
{
|
|
std::vector<uint32_t> h(cacheWords);
|
|
CUDA_CHECK(cudaMemcpy(h.data(), dCache, (size_t)cacheWords * 4, cudaMemcpyDeviceToHost));
|
|
uint64_t fnv = fnv1a64(h.data(), (size_t)cacheWords * 4);
|
|
bool ok = fnv == IGNEUM_CACHE_FNV64 && std::memcmp(h.data(), IGNEUM_CACHE_HEAD, 64) == 0 && std::memcmp(h.data() + cacheWords - 16, IGNEUM_CACHE_LAST, 64) == 0;
|
|
std::printf("check cache: %s (GPU fill %.0f ms, FNV %016llx vs pack %016llx)\n", ok ? "PASS" : "FAIL", fillMs, (unsigned long long)fnv, (unsigned long long)IGNEUM_CACHE_FNV64);
|
|
if (!ok) return 1;
|
|
}
|
|
// 2. dataset
|
|
const uint32_t words = 1u << IGNEUM_DATASET_LOG2;
|
|
const uint32_t mask = words - 1u;
|
|
uint32_t* dDs = nullptr;
|
|
CUDA_CHECK(cudaMalloc((void**)&dDs, (size_t)words * 4));
|
|
t0 = nowMs();
|
|
CUDA_CHECK(igneum_launch_build(dDs, dCache, words / 16u));
|
|
CUDA_CHECK(cudaDeviceSynchronize());
|
|
double buildMs = nowMs() - t0;
|
|
{
|
|
int bad = 0;
|
|
for (int k = 0; k < IGNEUM_DS_SAMPLES; ++k) {
|
|
uint32_t v = 0;
|
|
CUDA_CHECK(cudaMemcpy(&v, dDs + IGNEUM_DS_SAMPLE_INDEX[k], 4, cudaMemcpyDeviceToHost));
|
|
if (v != IGNEUM_DS_SAMPLE_VALUE[k]) ++bad;
|
|
}
|
|
uint32_t head[16]; CUDA_CHECK(cudaMemcpy(head, dDs, 64, cudaMemcpyDeviceToHost));
|
|
uint32_t last = 0; CUDA_CHECK(cudaMemcpy(&last, dDs + IGNEUM_DS_LAST_INDEX, 4, cudaMemcpyDeviceToHost));
|
|
bool ok = bad == 0 && std::memcmp(head, IGNEUM_DS_HEAD, 64) == 0 && last == IGNEUM_DS_LAST;
|
|
std::printf("check dataset: %s (GPU build %.0f ms, %d of %d samples wrong)\n", ok ? "PASS" : "FAIL", buildMs, bad, (int)IGNEUM_DS_SAMPLES);
|
|
if (!ok) return 1;
|
|
}
|
|
// 3. the pack's own vectors through the bench kernel (init words = seed words)
|
|
const uint32_t batch = 1u << batchLog2;
|
|
const size_t warps = batch / 32;
|
|
uint64_t* dOut = nullptr;
|
|
CUDA_CHECK(cudaMalloc((void**)&dOut, (size_t)batch * 8));
|
|
std::vector<uint64_t> hOut(batch);
|
|
{
|
|
int bad = 0;
|
|
for (int w = 0; w < IGNEUM_VEC_WARPS; ++w) {
|
|
CUDA_CHECK(igneum_launch_hash(dDs, dOut, IGNEUM_VEC_BASE[w], mask, 32u, 1u));
|
|
CUDA_CHECK(cudaDeviceSynchronize());
|
|
CUDA_CHECK(cudaMemcpy(hOut.data(), dOut, 32 * 8, cudaMemcpyDeviceToHost));
|
|
for (int l = 0; l < 32; ++l) if (hOut[l] != IGNEUM_VEC_OUT[w][l]) ++bad;
|
|
}
|
|
std::printf("check vectors: %s (%d of 96 lanes wrong)\n", bad == 0 ? "PASS" : "FAIL", bad);
|
|
if (bad) return 1;
|
|
}
|
|
// 4. the bound hash against igneum-pow's reference (nonce_hi 0, lane nonces 0..63)
|
|
IgneumInitWords iw0 = blockInit(ph.data(), 0u);
|
|
std::vector<uint64_t> honest64(64);
|
|
{
|
|
CUDA_CHECK(igneum_launch_hash_bound(dDs, dOut, 0u, mask, iw0, 64u, 1u));
|
|
CUDA_CHECK(cudaDeviceSynchronize());
|
|
CUDA_CHECK(cudaMemcpy(honest64.data(), dOut, 64 * 8, cudaMemcpyDeviceToHost));
|
|
int bad = 0, n = 0;
|
|
std::ifstream rf(refPath);
|
|
std::string line;
|
|
while (std::getline(rf, line)) {
|
|
if (line.empty() || line[0] == '#') continue;
|
|
unsigned long long nonce = 0, hash = 0;
|
|
if (std::sscanf(line.c_str(), "%llu %llx", &nonce, &hash) != 2 || nonce >= 64) continue;
|
|
++n;
|
|
if (honest64[nonce] != hash) ++bad;
|
|
}
|
|
std::printf("check bound vs igneum-pow: %s (%d reference lines, %d wrong; init %08x %08x ...)\n", (n == 64 && bad == 0) ? "PASS" : "FAIL", n, bad, iw0.w[0], iw0.w[1]);
|
|
if (n != 64 || bad) return 1;
|
|
}
|
|
// 5. the per-warp kernel on an all-equal table reproduces the honest kernel (known-pass); the random table's
|
|
// second warp differs (known-fail); the forced kernels differ (known-fail)
|
|
IgneumInitWords* dTbl = nullptr;
|
|
CUDA_CHECK(cudaMalloc((void**)&dTbl, warps * sizeof(IgneumInitWords)));
|
|
{
|
|
std::vector<IgneumInitWords> same(2, iw0);
|
|
CUDA_CHECK(cudaMemcpy(dTbl, same.data(), 2 * sizeof(IgneumInitWords), cudaMemcpyHostToDevice));
|
|
CUDA_CHECK(f9_launch_perwarp(dDs, dOut, 0u, mask, dTbl, 64u, 1u));
|
|
CUDA_CHECK(cudaDeviceSynchronize());
|
|
CUDA_CHECK(cudaMemcpy(hOut.data(), dOut, 64 * 8, cudaMemcpyDeviceToHost));
|
|
int diff = 0; for (int l = 0; l < 64; ++l) diff += hOut[l] != honest64[l];
|
|
std::printf("check perwarp all-equal table == honest: %s (%d of 64 differ)\n", diff == 0 ? "PASS" : "FAIL", diff);
|
|
if (diff) return 1;
|
|
std::vector<IgneumInitWords> two = { iw0, blockInit(ph.data(), 1u) };
|
|
CUDA_CHECK(cudaMemcpy(dTbl, two.data(), 2 * sizeof(IgneumInitWords), cudaMemcpyHostToDevice));
|
|
CUDA_CHECK(f9_launch_perwarp(dDs, dOut, 0u, mask, dTbl, 64u, 1u));
|
|
CUDA_CHECK(cudaDeviceSynchronize());
|
|
CUDA_CHECK(cudaMemcpy(hOut.data(), dOut, 64 * 8, cudaMemcpyDeviceToHost));
|
|
int d0 = 0, d1 = 0; for (int l = 0; l < 32; ++l) { d0 += hOut[l] != honest64[l]; d1 += hOut[32 + l] != honest64[32 + l]; }
|
|
std::printf("check perwarp two-init table: warp 0 equal (%d differ), warp 1 differs (%d of 32): %s (known-fail fires)\n", d0, d1, (d0 == 0 && d1 == 32) ? "PASS" : "FAIL");
|
|
if (d0 || d1 != 32) return 1;
|
|
CUDA_CHECK(f9_launch_forced4(dDs, dOut, 0u, mask, iw0, 64u, 1u));
|
|
CUDA_CHECK(cudaDeviceSynchronize());
|
|
CUDA_CHECK(cudaMemcpy(hOut.data(), dOut, 64 * 8, cudaMemcpyDeviceToHost));
|
|
int f4 = 0; for (int l = 0; l < 64; ++l) f4 += hOut[l] != honest64[l];
|
|
CUDA_CHECK(f9_launch_forcedall(dDs, dOut, 0u, mask, iw0, 64u, 1u));
|
|
CUDA_CHECK(cudaDeviceSynchronize());
|
|
CUDA_CHECK(cudaMemcpy(hOut.data(), dOut, 64 * 8, cudaMemcpyDeviceToHost));
|
|
int fa = 0; for (int l = 0; l < 64; ++l) fa += hOut[l] != honest64[l];
|
|
CUDA_CHECK(f9_launch_forced1(dDs, dOut, 0u, mask, iw0, 64u, 1u));
|
|
CUDA_CHECK(cudaDeviceSynchronize());
|
|
CUDA_CHECK(cudaMemcpy(hOut.data(), dOut, 64 * 8, cudaMemcpyDeviceToHost));
|
|
int f1 = 0; for (int l = 0; l < 64; ++l) f1 += hOut[l] != honest64[l];
|
|
CUDA_CHECK(f9_launch_pair(dDs, dOut, 0u, mask, iw0, 64u, 1u));
|
|
CUDA_CHECK(cudaDeviceSynchronize());
|
|
CUDA_CHECK(cudaMemcpy(hOut.data(), dOut, 64 * 8, cudaMemcpyDeviceToHost));
|
|
int pr = 0; for (int l = 0; l < 64; ++l) pr += hOut[l] != honest64[l];
|
|
std::printf("check forced kernels change the hash (known-fail fires): forced4 %d of 64 differ, forcedall %d, forced1 %d, pair %d: %s\n", f4, fa, f1, pr, (f4 >= 60 && fa >= 60 && f1 >= 60 && pr >= 2) ? "PASS" : "FAIL");
|
|
if (f4 < 60 || fa < 60 || f1 < 60 || pr < 2) return 1;
|
|
}
|
|
// tables
|
|
std::vector<IgneumInitWords> tblRandom = loadTable(tRandom, warps), tblK10 = loadTable(tK10, warps), tblK14 = loadTable(tK14, warps);
|
|
{
|
|
// the random table's warp 0 is nonce_hi 0: equal to iw0 by construction
|
|
bool ok = std::memcmp(&tblRandom[0], &iw0, sizeof iw0) == 0;
|
|
std::printf("check table-random warp 0 == init(nonce_hi 0): %s\n", ok ? "PASS" : "FAIL");
|
|
if (!ok) return 1;
|
|
}
|
|
std::printf("timed rounds: %d rounds x %d variants, %d s each, batch 2^%d nonces, %d warps per block\n", rounds, (int)NVARIANTS, seconds, batchLog2, blockWarps);
|
|
|
|
cudaEvent_t e0, e1;
|
|
CUDA_CHECK(cudaEventCreate(&e0));
|
|
CUDA_CHECK(cudaEventCreate(&e1));
|
|
std::vector<std::vector<double>> rate(NVARIANTS);
|
|
for (int r = 0; r < rounds; ++r) {
|
|
for (int v = 0; v < NVARIANTS; ++v) {
|
|
const std::vector<IgneumInitWords>* tbl = v == PERWARP_RANDOM ? &tblRandom : v == PERWARP_K10 ? &tblK10 : v == PERWARP_K14 ? &tblK14 : nullptr;
|
|
if (tbl) CUDA_CHECK(cudaMemcpy(dTbl, tbl->data(), warps * sizeof(IgneumInitWords), cudaMemcpyHostToDevice));
|
|
auto launch = [&](uint32_t nonceHi) -> cudaError_t {
|
|
IgneumInitWords iw = blockInit(ph.data(), nonceHi);
|
|
switch (v) {
|
|
case HONEST: return igneum_launch_hash_bound(dDs, dOut, 0u, mask, iw, batch, (uint32_t)blockWarps);
|
|
case PERWARP_RANDOM: case PERWARP_K10: case PERWARP_K14: return f9_launch_perwarp(dDs, dOut, 0u, mask, dTbl, batch, (uint32_t)blockWarps);
|
|
case FORCED4: return f9_launch_forced4(dDs, dOut, 0u, mask, iw, batch, (uint32_t)blockWarps);
|
|
case FORCED1: return f9_launch_forced1(dDs, dOut, 0u, mask, iw, batch, (uint32_t)blockWarps);
|
|
case PAIR: return f9_launch_pair(dDs, dOut, 0u, mask, iw, batch, (uint32_t)blockWarps);
|
|
default: return f9_launch_forcedall(dDs, dOut, 0u, mask, iw, batch, (uint32_t)blockWarps);
|
|
}
|
|
};
|
|
// warm-up
|
|
CUDA_CHECK(launch(0xffffffffu));
|
|
CUDA_CHECK(cudaDeviceSynchronize());
|
|
long long start = epochMs();
|
|
double wall0 = nowMs();
|
|
double gpuMs = 0;
|
|
uint64_t hashes = 0;
|
|
uint32_t d = 0;
|
|
while (nowMs() - wall0 < seconds * 1000.0) {
|
|
CUDA_CHECK(cudaEventRecord(e0));
|
|
CUDA_CHECK(launch(d));
|
|
CUDA_CHECK(cudaEventRecord(e1));
|
|
CUDA_CHECK(cudaEventSynchronize(e1));
|
|
float ms = 0; CUDA_CHECK(cudaEventElapsedTime(&ms, e0, e1));
|
|
gpuMs += ms;
|
|
hashes += batch;
|
|
++d;
|
|
}
|
|
long long end = epochMs();
|
|
double mhs = hashes / (gpuMs * 1e3);
|
|
rate[v].push_back(mhs);
|
|
std::printf("phase %s round %d start_ms %lld end_ms %lld dispatches %u hashes %llu gpu_ms %.1f MHs %.4f\n", NAMES[v], r, start, end, d, (unsigned long long)hashes, gpuMs, mhs);
|
|
std::fflush(stdout);
|
|
}
|
|
}
|
|
std::printf("summary (MH/s per variant over %d rounds: mean, sd, relative to honest)\n", rounds);
|
|
double hm = 0; for (double x : rate[HONEST]) hm += x; hm /= rate[HONEST].size();
|
|
for (int v = 0; v < NVARIANTS; ++v) {
|
|
double m = 0, s = 0; for (double x : rate[v]) m += x; m /= rate[v].size();
|
|
for (double x : rate[v]) s += (x - m) * (x - m); s = rate[v].size() > 1 ? std::sqrt(s / (rate[v].size() - 1)) : 0;
|
|
std::printf("variant %-15s mean %.4f sd %.4f rel %+.3f%%\n", NAMES[v], m, s, 100.0 * (m - hm) / hm);
|
|
}
|
|
cudaFree(dTbl); cudaFree(dOut); cudaFree(dDs); cudaFree(dCache);
|
|
return 0;
|
|
}
|