From 18166dbc46c1e0c2cb1e17b37e875ae153fcf458 Mon Sep 17 00:00:00 2001 From: Dileep-Kumar-5 Date: Wed, 5 Aug 2026 22:56:39 +0530 Subject: [PATCH 1/7] feat: widen kangaroo distance field to 256 bits (125 -> 253 bit intervals) The 125-bit interval limit was a struct field width, not a mathematical one. ENTRY.d packed the travelled distance into 128 bits as "b127=sign b126=kangaroo type, b125..b0 distance", leaving 126 bits of magnitude. Distances span the interval width, so anything wider than 125 bits overflowed. It overflowed silently. Convert() masked the excess away with & 0x3FFFFFFFFFFFFFFF, stored a valid-looking entry, and on collision produced a wrong private key with no error, no warning and no crash. Measured on the real puzzle #140 range (2^139): stock JLP stores 376,747 corrupt DPs out of 376,836 (-wcheck reports 0.024% OK) while reporting normal progress throughout. Distance is now 256 bits, flags in the top two bits, 254 bits of magnitude -> intervals up to 253 bits. Over-range distances abort instead of truncating. Storage and host side: - int256_t; ENTRY 32 -> 48 bytes; ENTRY_SIZE replaces hard-coded 32/16 - sameDist() compares all four words; the old check compared two of four and would mis-flag a real collision as a duplicate - work file magics bumped: old files are rejected, not misparsed GPU: - kernel dist[GPU_GRP_SIZE] 2 -> 4 words, KSIZE 10 -> 12 (11 -> 13 sym) - new Add256 carries through all four words - jD[NB_JUMP] 2 -> 4 words; SetParams uploads the full jump table - ITEM_SIZE 56 -> 72; OutputDP writes 8 distance words, kIdx at +17 - fixes a latent out-of-bounds write: ModNeg256Order(dist[g]) wrote r[0..3] into a 2-word dist[g], clobbering the next kangaroo. Only dormant because USE_SYMMETRY ships commented out. Network: - DP packet 40 -> 56 bytes, all four distance words sent - kangaroo blocks 16 -> 32 bytes, checksums cover all four words - SERVER_HEADER bumped so old peers fail the handshake The now-unreachable 126-bit Add/Convert/CalcDistAndType/Widen overloads are removed. Verified: CPU build (MSVC) solves the 56-bit sample key correctly, and on the 2^139 puzzle #140 range stores 408,118 DPs at 100.000% -wcheck. NOT verified: the GPU path has never been through nvcc, and client/ server mode compiles but was not exercised at runtime. Costs: DP tables +50% RAM, device kangaroo memory +20%, DP packets +40%, kangaroo transfers +100%. Wider dist raises local-memory pressure in the kernel; expect to lower -g. Co-Authored-By: Claude Opus 5 --- .gitignore | 9 +++++ Backup.cpp | 10 +++--- Check.cpp | 2 +- GPU/GPUCompute.h | 4 +-- GPU/GPUEngine.cu | 20 +++++++---- GPU/GPUEngine.h | 9 +++-- GPU/GPUMath.h | 45 ++++++++++++++++++------ HashTable.cpp | 89 ++++++++++++++++++++++++++++++------------------ HashTable.h | 38 ++++++++++++++++++--- Kangaroo.cpp | 2 +- Kangaroo.h | 19 ++++++----- Network.cpp | 52 +++++++++++++++++----------- build_cpu.bat | 14 ++++++++ build_gpu.bat | 47 +++++++++++++++++++++++++ test140.txt | 3 ++ 15 files changed, 267 insertions(+), 96 deletions(-) create mode 100644 build_cpu.bat create mode 100644 build_gpu.bat create mode 100644 test140.txt diff --git a/.gitignore b/.gitignore index b293acf6..a89636b9 100644 --- a/.gitignore +++ b/.gitignore @@ -12,3 +12,12 @@ CMakeLists.txt kangaroo cmake-build-debug/ *.prf + +# Windows MSVC/nvcc build output (build_cpu.bat / build_gpu.bat) +objgpu/ +*.obj +*.exe + +# run artifacts +*.kcp +*.log diff --git a/Backup.cpp b/Backup.cpp index e90d74f4..71db105b 100644 --- a/Backup.cpp +++ b/Backup.cpp @@ -230,7 +230,7 @@ void Kangaroo::FetchWalks(uint64_t nbWalk,Int *x,Int *y,Int *d) { } -void Kangaroo::FetchWalks(uint64_t nbWalk,std::vector& kangs,Int* x,Int* y,Int* d) { +void Kangaroo::FetchWalks(uint64_t nbWalk,std::vector& kangs,Int* x,Int* y,Int* d) { uint64_t n = 0; @@ -293,7 +293,7 @@ void Kangaroo::FectchKangaroos(TH_PARAM *threads) { double sFetch = Timer::get_tick(); // From server - vector kangs; + vector kangs; if(saveKangarooByServer) { ::printf("FectchKangaroosFromServer"); if(!GetKangaroosFromServer(workFile,kangs)) @@ -492,14 +492,14 @@ void Kangaroo::SaveWork(uint64_t totalCount,double totalTime,TH_PARAM *threads,i if(saveKangarooByServer) { ::printf("\nSaveWork (Kangaroo->Server): %s",fileName.c_str()); - vector kangs; + vector kangs; for(int i = 0; i < nbThread; i++) totalWalk += threads[i].nbKangaroo; kangs.reserve(totalWalk); for(int i = 0; i < nbThread; i++) { int128_t X; - int128_t D; + int256_t D; uint64_t h; for(uint64_t n = 0; n < threads[i].nbKangaroo; n++) { HashTable::Convert(&threads[i].px[n],&threads[i].distance[n],n%2,&h,&X,&D); @@ -507,7 +507,7 @@ void Kangaroo::SaveWork(uint64_t totalCount,double totalTime,TH_PARAM *threads,i } } SendKangaroosToServer(fileName,kangs); - size = kangs.size()*16 + 16; + size = kangs.size()*32 + 16; goto end; } else { diff --git a/Check.cpp b/Check.cpp index 8eeb4f05..cbfae3e8 100644 --- a/Check.cpp +++ b/Check.cpp @@ -57,7 +57,7 @@ uint32_t Kangaroo::CheckHash(uint32_t h,uint32_t nbItem,HashTable* hT,FILE* f) { items = (ENTRY*)malloc(nbItem * sizeof(ENTRY)); for(uint32_t i = 0; i < nbItem; i++) { - ::fread(items+i,32,1,f); + ::fread(items+i,ENTRY_SIZE,1,f); e = items + i; Int dist; uint32_t kType; diff --git a/GPU/GPUCompute.h b/GPU/GPUCompute.h index 95a358a2..ff03fe0a 100644 --- a/GPU/GPUCompute.h +++ b/GPU/GPUCompute.h @@ -23,7 +23,7 @@ __device__ void ComputeKangaroos(uint64_t *kangaroos,uint32_t maxFound,uint32_t uint64_t px[GPU_GRP_SIZE][4]; uint64_t py[GPU_GRP_SIZE][4]; - uint64_t dist[GPU_GRP_SIZE][2]; + uint64_t dist[GPU_GRP_SIZE][4]; #ifdef USE_SYMMETRY uint64_t lastJump[GPU_GRP_SIZE]; #endif @@ -86,7 +86,7 @@ __device__ void ComputeKangaroos(uint64_t *kangaroos,uint32_t maxFound,uint32_t Load256(px[g],rx); Load256(py[g],ry); - Add128(dist[g],jD[jmp]); + Add256(dist[g],jD[jmp]); #ifdef USE_SYMMETRY if(ModPositive256(py[g])) diff --git a/GPU/GPUEngine.cu b/GPU/GPUEngine.cu index 86be42b4..bfdc49e8 100644 --- a/GPU/GPUEngine.cu +++ b/GPU/GPUEngine.cu @@ -409,10 +409,12 @@ void GPUEngine::SetKangaroos(Int *px,Int *py,Int *d) { if(idx % 2 == WILD) dOff.ModAddK1order(&wildOffset); inputKangarooPinned[g * strideSize + t + 8 * nbThreadPerGroup] = dOff.bits64[0]; inputKangarooPinned[g * strideSize + t + 9 * nbThreadPerGroup] = dOff.bits64[1]; + inputKangarooPinned[g * strideSize + t + 10 * nbThreadPerGroup] = dOff.bits64[2]; + inputKangarooPinned[g * strideSize + t + 11 * nbThreadPerGroup] = dOff.bits64[3]; #ifdef USE_SYMMETRY // Last jump - inputKangarooPinned[t + 10 * nbThreadPerGroup] = (uint64_t)NB_JUMP; + inputKangarooPinned[t + 12 * nbThreadPerGroup] = (uint64_t)NB_JUMP; #endif idx++; @@ -474,6 +476,8 @@ void GPUEngine::GetKangaroos(Int *px,Int *py,Int *d) { dOff.SetInt32(0); dOff.bits64[0] = inputKangarooPinned[g * strideSize + t + 8 * nbThreadPerGroup]; dOff.bits64[1] = inputKangarooPinned[g * strideSize + t + 9 * nbThreadPerGroup]; + dOff.bits64[2] = inputKangarooPinned[g * strideSize + t + 10 * nbThreadPerGroup]; + dOff.bits64[3] = inputKangarooPinned[g * strideSize + t + 11 * nbThreadPerGroup]; if(idx % 2 == WILD) dOff.ModSubK1order(&wildOffset); d[idx].Set(&dOff); @@ -528,6 +532,10 @@ void GPUEngine::SetKangaroo(uint64_t kIdx,Int *px,Int *py,Int *d) { cudaMemcpy(inputKangaroo + (b * blockSize + g * strideSize + t + 8 * nbThreadPerGroup),inputKangarooPinned,8,cudaMemcpyHostToDevice); inputKangarooPinned[0] = dOff.bits64[1]; cudaMemcpy(inputKangaroo + (b * blockSize + g * strideSize + t + 9 * nbThreadPerGroup),inputKangarooPinned,8,cudaMemcpyHostToDevice); + inputKangarooPinned[0] = dOff.bits64[2]; + cudaMemcpy(inputKangaroo + (b * blockSize + g * strideSize + t + 10 * nbThreadPerGroup),inputKangarooPinned,8,cudaMemcpyHostToDevice); + inputKangarooPinned[0] = dOff.bits64[3]; + cudaMemcpy(inputKangaroo + (b * blockSize + g * strideSize + t + 11 * nbThreadPerGroup),inputKangarooPinned,8,cudaMemcpyHostToDevice); #ifdef USE_SYMMETRY // Last jump @@ -561,8 +569,8 @@ void GPUEngine::SetParams(uint64_t dpMask,Int *distance,Int *px,Int *py) { this->dpMask = dpMask; for(int i=0;i< NB_JUMP;i++) - memcpy(jumpPinned + 2*i,distance[i].bits64,16); - cudaMemcpyToSymbol(jD,jumpPinned,jumpSize/2); + memcpy(jumpPinned + 4*i,distance[i].bits64,32); + cudaMemcpyToSymbol(jD,jumpPinned,jumpSize); cudaError_t err = cudaGetLastError(); if(err != cudaSuccess) { printf("GPUEngine: SetParams: Failed to copy to constant memory: %s\n",cudaGetErrorString(err)); @@ -654,7 +662,7 @@ bool GPUEngine::Launch(std::vector &hashFound,bool spinWait) { uint32_t *itemPtr = outputItemPinned + (i*ITEM_SIZE32 + 1); ITEM it; - it.kIdx = *((uint64_t*)(itemPtr + 12)); + it.kIdx = *((uint64_t*)(itemPtr + 16)); uint64_t *x = (uint64_t *)itemPtr; it.x.bits64[0] = x[0]; @@ -666,8 +674,8 @@ bool GPUEngine::Launch(std::vector &hashFound,bool spinWait) { uint64_t *d = (uint64_t *)(itemPtr + 8); it.d.bits64[0] = d[0]; it.d.bits64[1] = d[1]; - it.d.bits64[2] = 0; - it.d.bits64[3] = 0; + it.d.bits64[2] = d[2]; + it.d.bits64[3] = d[3]; it.d.bits64[4] = 0; if(it.kIdx % 2 == WILD) it.d.ModSubK1order(&wildOffset); diff --git a/GPU/GPUEngine.h b/GPU/GPUEngine.h index 8f61099e..fe1d00a7 100644 --- a/GPU/GPUEngine.h +++ b/GPU/GPUEngine.h @@ -22,13 +22,16 @@ #include "../Constants.h" #include "../SECPK1/SECP256k1.h" +// Words per kangaroo in device memory: px[4] + py[4] + dist[4] (+ lastJump). +// dist was 2 words (126bit distance, 125bit interval cap); it is 4 words now. #ifdef USE_SYMMETRY -#define KSIZE 11 +#define KSIZE 13 #else -#define KSIZE 10 +#define KSIZE 12 #endif -#define ITEM_SIZE 56 +// x[8] + d[8] + kIdx[2], in uint32 +#define ITEM_SIZE 72 #define ITEM_SIZE32 (ITEM_SIZE/4) typedef struct { diff --git a/GPU/GPUMath.h b/GPU/GPUMath.h index b67e0055..ae008c31 100644 --- a/GPU/GPUMath.h +++ b/GPU/GPUMath.h @@ -48,7 +48,7 @@ #define MADDS(r,a,b,c) asm volatile ("madc.hi.s64 %0, %1, %2, %3;" : "=l"(r) : "l"(a), "l"(b), "l"(c)); // Jump distance -__device__ __constant__ uint64_t jD[NB_JUMP][2]; +__device__ __constant__ uint64_t jD[NB_JUMP][4]; // jump points __device__ __constant__ uint64_t jPx[NB_JUMP][4]; __device__ __constant__ uint64_t jPy[NB_JUMP][4]; @@ -122,6 +122,17 @@ __device__ __constant__ uint64_t _O[] = { 0xBFD25E8CD0364141ULL,0xBAAEDCE6AF48A0 // --------------------------------------------------------------------------------------- +// 256bit accumulate, for the kangaroo travelled distance. Carry has to run +// through all four words -- truncating at 128 bits is what capped the search +// interval at 125 bits and silently corrupted the recovered key beyond it. +#define Add256(r,a) { \ + UADDO1((r)[0], (a)[0]); \ + UADDC1((r)[1], (a)[1]); \ + UADDC1((r)[2], (a)[2]); \ + UADD1((r)[3], (a)[3]);} + +// --------------------------------------------------------------------------------------- + #define Neg(r) {\ USUBO(r[0],0ULL,r[0]); \ USUBC(r[1],0ULL,r[1]); \ @@ -183,16 +194,20 @@ out[pos*ITEM_SIZE32 + 9] = ((uint32_t *)d)[0]; \ out[pos*ITEM_SIZE32 + 10] = ((uint32_t *)d)[1]; \ out[pos*ITEM_SIZE32 + 11] = ((uint32_t *)d)[2]; \ out[pos*ITEM_SIZE32 + 12] = ((uint32_t *)d)[3]; \ -out[pos*ITEM_SIZE32 + 13] = ((uint32_t *)idx)[0]; \ -out[pos*ITEM_SIZE32 + 14] = ((uint32_t *)idx)[1]; \ +out[pos*ITEM_SIZE32 + 13] = ((uint32_t *)d)[4]; \ +out[pos*ITEM_SIZE32 + 14] = ((uint32_t *)d)[5]; \ +out[pos*ITEM_SIZE32 + 15] = ((uint32_t *)d)[6]; \ +out[pos*ITEM_SIZE32 + 16] = ((uint32_t *)d)[7]; \ +out[pos*ITEM_SIZE32 + 17] = ((uint32_t *)idx)[0]; \ +out[pos*ITEM_SIZE32 + 18] = ((uint32_t *)idx)[1]; \ } // --------------------------------------------------------------------------------------- #ifdef USE_SYMMETRY -__device__ void LoadKangaroos(uint64_t *a,uint64_t px[GPU_GRP_SIZE][4],uint64_t py[GPU_GRP_SIZE][4],uint64_t dist[GPU_GRP_SIZE][2],uint64_t *jumps) { +__device__ void LoadKangaroos(uint64_t *a,uint64_t px[GPU_GRP_SIZE][4],uint64_t py[GPU_GRP_SIZE][4],uint64_t dist[GPU_GRP_SIZE][4],uint64_t *jumps) { #else -__device__ void LoadKangaroos(uint64_t * a,uint64_t px[GPU_GRP_SIZE][4],uint64_t py[GPU_GRP_SIZE][4],uint64_t dist[GPU_GRP_SIZE][2]) { +__device__ void LoadKangaroos(uint64_t * a,uint64_t px[GPU_GRP_SIZE][4],uint64_t py[GPU_GRP_SIZE][4],uint64_t dist[GPU_GRP_SIZE][4]) { #endif __syncthreads(); @@ -216,15 +231,17 @@ __device__ void LoadKangaroos(uint64_t * a,uint64_t px[GPU_GRP_SIZE][4],uint64_t d64[0] = (a)[IDX + 8 * blockDim.x + stride]; d64[1] = (a)[IDX + 9 * blockDim.x + stride]; + d64[2] = (a)[IDX + 10 * blockDim.x + stride]; + d64[3] = (a)[IDX + 11 * blockDim.x + stride]; #ifdef USE_SYMMETRY - jumps[g] = (a)[IDX + 10 * blockDim.x + stride]; + jumps[g] = (a)[IDX + 12 * blockDim.x + stride]; #endif } } -__device__ void LoadDists(uint64_t* a,uint64_t dist[GPU_GRP_SIZE][2]) { +__device__ void LoadDists(uint64_t* a,uint64_t dist[GPU_GRP_SIZE][4]) { __syncthreads(); @@ -235,6 +252,8 @@ __device__ void LoadDists(uint64_t* a,uint64_t dist[GPU_GRP_SIZE][2]) { d64[0] = (a)[IDX + 8 * blockDim.x + stride]; d64[1] = (a)[IDX + 9 * blockDim.x + stride]; + d64[2] = (a)[IDX + 10 * blockDim.x + stride]; + d64[3] = (a)[IDX + 11 * blockDim.x + stride]; } @@ -271,9 +290,9 @@ __device__ void LoadKangaroo(uint64_t* a,uint32_t stride,uint64_t px[4]) { // --------------------------------------------------------------------------------------- #ifdef USE_SYMMETRY -__device__ void StoreKangaroos(uint64_t *a,uint64_t px[GPU_GRP_SIZE][4],uint64_t py[GPU_GRP_SIZE][4],uint64_t dist[GPU_GRP_SIZE][2],uint64_t *jumps) { +__device__ void StoreKangaroos(uint64_t *a,uint64_t px[GPU_GRP_SIZE][4],uint64_t py[GPU_GRP_SIZE][4],uint64_t dist[GPU_GRP_SIZE][4],uint64_t *jumps) { #else -__device__ void StoreKangaroos(uint64_t * a,uint64_t px[GPU_GRP_SIZE][4],uint64_t py[GPU_GRP_SIZE][4],uint64_t dist[GPU_GRP_SIZE][2]) { +__device__ void StoreKangaroos(uint64_t * a,uint64_t px[GPU_GRP_SIZE][4],uint64_t py[GPU_GRP_SIZE][4],uint64_t dist[GPU_GRP_SIZE][4]) { #endif __syncthreads(); @@ -296,9 +315,11 @@ __device__ void StoreKangaroos(uint64_t * a,uint64_t px[GPU_GRP_SIZE][4],uint64_ (a)[IDX + 8 * blockDim.x + stride] = d64[0]; (a)[IDX + 9 * blockDim.x + stride] = d64[1]; + (a)[IDX + 10 * blockDim.x + stride] = d64[2]; + (a)[IDX + 11 * blockDim.x + stride] = d64[3]; #ifdef USE_SYMMETRY - (a)[IDX + 10 * blockDim.x + stride] = jumps[g]; + (a)[IDX + 12 * blockDim.x + stride] = jumps[g]; #endif } @@ -321,7 +342,7 @@ __device__ void StoreKangaroo(uint64_t* a,uint32_t stride,uint64_t px[4],uint64_ } -__device__ void StoreDists(uint64_t* a,uint64_t dist[GPU_GRP_SIZE][2]) { +__device__ void StoreDists(uint64_t* a,uint64_t dist[GPU_GRP_SIZE][4]) { __syncthreads(); @@ -331,6 +352,8 @@ __device__ void StoreDists(uint64_t* a,uint64_t dist[GPU_GRP_SIZE][2]) { (a)[IDX + 8 * blockDim.x + stride] = d64[0]; (a)[IDX + 9 * blockDim.x + stride] = d64[1]; + (a)[IDX + 10 * blockDim.x + stride] = d64[2]; + (a)[IDX + 11 * blockDim.x + stride] = d64[3]; } diff --git a/HashTable.cpp b/HashTable.cpp index 20ab658c..51c97969 100644 --- a/HashTable.cpp +++ b/HashTable.cpp @@ -17,6 +17,7 @@ #include "HashTable.h" #include +#include #include #ifndef WIN64 #include @@ -54,17 +55,26 @@ uint64_t HashTable::GetNbItem() { } -ENTRY *HashTable::CreateEntry(int128_t *x,int128_t *d) { +ENTRY *HashTable::CreateEntry(int128_t *x,int256_t *d) { ENTRY *e = (ENTRY *)malloc(sizeof(ENTRY)); e->x.i64[0] = x->i64[0]; e->x.i64[1] = x->i64[1]; e->d.i64[0] = d->i64[0]; e->d.i64[1] = d->i64[1]; + e->d.i64[2] = d->i64[2]; + e->d.i64[3] = d->i64[3]; return e; } +bool HashTable::sameDist(int256_t *a,int256_t *b) { + + return (a->i64[0] == b->i64[0]) && (a->i64[1] == b->i64[1]) && + (a->i64[2] == b->i64[2]) && (a->i64[3] == b->i64[3]); + +} + #define ADD_ENTRY(entry) { \ /* Shift the end of the index table */ \ for (int i = E[h].nbItem; i > st; i--) \ @@ -72,7 +82,7 @@ ENTRY *HashTable::CreateEntry(int128_t *x,int128_t *d) { E[h].items[st] = entry; \ E[h].nbItem++;} -void HashTable::Convert(Int *x,Int *d,uint32_t type,uint64_t *h,int128_t *X,int128_t *D) { +void HashTable::Convert(Int *x,Int *d,uint32_t type,uint64_t *h,int128_t *X,int256_t *D) { uint64_t sign = 0; uint64_t type64 = (uint64_t)type << 62; @@ -80,28 +90,37 @@ void HashTable::Convert(Int *x,Int *d,uint32_t type,uint64_t *h,int128_t *X,int1 X->i64[0] = x->bits64[0]; X->i64[1] = x->bits64[1]; - // Probability of failure (1/2^128) + Int N(d); + + // Distances are held mod n; a value in the upper half means "negative". + // Probability of failure (1/2^256) if(d->bits64[3] > 0x7FFFFFFFFFFFFFFFULL) { - Int N(d); N.ModNegK1order(); - D->i64[0] = N.bits64[0]; - D->i64[1] = N.bits64[1] & 0x3FFFFFFFFFFFFFFFULL; - sign = 1ULL << 63; - } else { - D->i64[0] = d->bits64[0]; - D->i64[1] = d->bits64[1] & 0x3FFFFFFFFFFFFFFFULL; + sign = DIST_SIGN_MASK; + } + + // The magnitude has to fit in b253..b0. The original code masked the excess + // away, which produced a valid-looking entry and, on collision, a WRONG + // private key with no error. Refuse loudly instead. + if(N.bits64[3] & 0xC000000000000000ULL) { + ::printf("\nHashTable::Convert: travelled distance exceeds %d bits.\n" + "Interval is too large for the DP entry format (max %d bits).\n" + "Aborting rather than storing a truncated distance.\n", + DIST_MAG_BITS,MAX_INTERVAL_BITS); + exit(-1); } - D->i64[1] |= sign; - D->i64[1] |= type64; + D->i64[0] = N.bits64[0]; + D->i64[1] = N.bits64[1]; + D->i64[2] = N.bits64[2]; + D->i64[3] = (N.bits64[3] & DIST_MAG_MASK) | sign | type64; *h = (x->bits64[2] & HASH_MASK); } - -#define AV1() if(pnb1) { ::fread(&e1,32,1,f1); pnb1--; } -#define AV2() if(pnb2) { ::fread(&e2,32,1,f2); pnb2--; } +#define AV1() if(pnb1) { ::fread(&e1,ENTRY_SIZE,1,f1); pnb1--; } +#define AV2() if(pnb2) { ::fread(&e2,ENTRY_SIZE,1,f2); pnb2--; } int HashTable::MergeH(uint32_t h,FILE* f1,FILE* f2,FILE* fd,uint32_t* nbDP,uint32_t *duplicate,Int* d1,uint32_t* k1,Int* d2,uint32_t* k2) { @@ -152,12 +171,12 @@ int HashTable::MergeH(uint32_t h,FILE* f1,FILE* f2,FILE* fd,uint32_t* nbDP,uint3 int comp = compare(&e1.x,&e2.x); if(comp < 0) { - memcpy(output+nbd,&e1,32); + memcpy(output+nbd,&e1,ENTRY_SIZE); nbd++; AV1(); nb1--; } else if (comp==0) { - if((e1.d.i64[0] == e2.d.i64[0]) && (e1.d.i64[1] == e2.d.i64[1])) { + if(sameDist(&e1.d,&e2.d)) { *duplicate = *duplicate + 1; } else { // Collision @@ -165,14 +184,14 @@ int HashTable::MergeH(uint32_t h,FILE* f1,FILE* f2,FILE* fd,uint32_t* nbDP,uint3 CalcDistAndType(e2.d,d2,k2); collisionFound = true; } - memcpy(output + nbd,&e1,32); + memcpy(output + nbd,&e1,ENTRY_SIZE); nbd++; AV1(); AV2(); nb1--; nb2--; } else { - memcpy(output + nbd,&e2,32); + memcpy(output + nbd,&e2,ENTRY_SIZE); nbd++; AV2(); nb2--; @@ -180,14 +199,14 @@ int HashTable::MergeH(uint32_t h,FILE* f1,FILE* f2,FILE* fd,uint32_t* nbDP,uint3 } else if( !end1 && end2 ) { - memcpy(output + nbd,&e1,32); + memcpy(output + nbd,&e1,ENTRY_SIZE); nbd++; AV1(); nb1--; } else if( end1 && !end2) { - memcpy(output + nbd,&e2,32); + memcpy(output + nbd,&e2,ENTRY_SIZE); nbd++; AV2(); nb2--; @@ -210,7 +229,7 @@ int HashTable::MergeH(uint32_t h,FILE* f1,FILE* f2,FILE* fd,uint32_t* nbDP,uint3 ::fwrite(&nbd,sizeof(uint32_t),1,fd); ::fwrite(&md,sizeof(uint32_t),1,fd); - ::fwrite(output,32,nbd,fd); + ::fwrite(output,ENTRY_SIZE,nbd,fd); free(output); *nbDP = nbd; @@ -221,7 +240,7 @@ int HashTable::MergeH(uint32_t h,FILE* f1,FILE* f2,FILE* fd,uint32_t* nbDP,uint3 int HashTable::Add(Int *x,Int *d,uint32_t type) { int128_t X; - int128_t D; + int256_t D; uint64_t h; Convert(x,d,type,&h,&X,&D); ENTRY* e = CreateEntry(&X,&D); @@ -239,22 +258,24 @@ void HashTable::ReAllocate(uint64_t h,uint32_t add) { } -int HashTable::Add(uint64_t h,int128_t *x,int128_t *d) { +int HashTable::Add(uint64_t h,int128_t *x,int256_t *d) { ENTRY *e = CreateEntry(x,d); return Add(h,e); } -void HashTable::CalcDistAndType(int128_t d,Int* kDist,uint32_t* kType) { +void HashTable::CalcDistAndType(int256_t d,Int* kDist,uint32_t* kType) { - *kType = (d.i64[1] & 0x4000000000000000ULL) != 0; - int sign = (d.i64[1] & 0x8000000000000000ULL) != 0; - d.i64[1] &= 0x3FFFFFFFFFFFFFFFULL; + *kType = (d.i64[3] & DIST_TYPE_MASK) != 0; + int sign = (d.i64[3] & DIST_SIGN_MASK) != 0; + d.i64[3] &= DIST_MAG_MASK; kDist->SetInt32(0); kDist->bits64[0] = d.i64[0]; kDist->bits64[1] = d.i64[1]; + kDist->bits64[2] = d.i64[2]; + kDist->bits64[3] = d.i64[3]; if(sign) kDist->ModNegK1order(); } @@ -287,7 +308,7 @@ int HashTable::Add(uint64_t h,ENTRY* e) { ed = mi - 1; } else if (comp==0) { - if((e->d.i64[0] == GET(h,mi)->d.i64[0]) && (e->d.i64[1] == GET(h,mi)->d.i64[1])) { + if(sameDist(&e->d,&GET(h,mi)->d)) { // Same point added 2 times or collision in same herd ! return ADD_DUPLICATE; } @@ -381,8 +402,8 @@ void HashTable::SaveTable(FILE* f,uint32_t from,uint32_t to,bool printPoint) { fwrite(&E[h].nbItem,sizeof(uint32_t),1,f); fwrite(&E[h].maxItem,sizeof(uint32_t),1,f); for(uint32_t i = 0; i < E[h].nbItem; i++) { - fwrite(&(E[h].items[i]->x),16,1,f); - fwrite(&(E[h].items[i]->d),16,1,f); + fwrite(&(E[h].items[i]->x),sizeof(int128_t),1,f); + fwrite(&(E[h].items[i]->d),sizeof(int256_t),1,f); if(printPoint) { pointPrint++; if(pointPrint > point) { @@ -425,7 +446,7 @@ void HashTable::SeekNbItem(FILE* f,uint32_t from,uint32_t to) { fread(&E[h].nbItem,sizeof(uint32_t),1,f); fread(&E[h].maxItem,sizeof(uint32_t),1,f); - uint64_t hSize = 32ULL * E[h].nbItem; + uint64_t hSize = (uint64_t)ENTRY_SIZE * E[h].nbItem; #ifdef WIN64 _fseeki64(f,hSize,SEEK_CUR); #else @@ -451,8 +472,8 @@ void HashTable::LoadTable(FILE* f,uint32_t from,uint32_t to) { for(uint32_t i = 0; i < E[h].nbItem; i++) { ENTRY* e = (ENTRY*)malloc(sizeof(ENTRY)); - fread(&(e->x),16,1,f); - fread(&(e->d),16,1,f); + fread(&(e->x),sizeof(int128_t),1,f); + fread(&(e->d),sizeof(int256_t),1,f); E[h].items[i] = e; } diff --git a/HashTable.h b/HashTable.h index e59d898d..9c1e4d1e 100644 --- a/HashTable.h +++ b/HashTable.h @@ -44,17 +44,44 @@ union int128_s { typedef union int128_s int128_t; +union int256_s { + + uint8_t i8[32]; + uint16_t i16[16]; + uint32_t i32[8]; + uint64_t i64[4]; + +}; + +typedef union int256_s int256_t; + #define safe_free(x) if(x) {free(x);x=NULL;} +// Distance field width. The sign and kangaroo-type flags live in the top two +// bits, exactly as in the original 128bit layout, so the magnitude gets +// b253..b0. That is 254 bits for a search that needs at most ~160, i.e. plenty +// of headroom -- moving the flags out to a separate byte would only pad ENTRY +// from 48 to 56 bytes and grow every DP table by 17% for nothing. +#define DIST_MAG_BITS 254 +#define MAX_INTERVAL_BITS 253 +#define DIST_SIGN_MASK 0x8000000000000000ULL // b255, in i64[3] +#define DIST_TYPE_MASK 0x4000000000000000ULL // b254, in i64[3] +#define DIST_MAG_MASK 0x3FFFFFFFFFFFFFFFULL // b253..b192, in i64[3] + // We store only 128 (+18) bit a the x value which give a probabilty a wrong collision after 2^73 entries typedef struct { int128_t x; // Poisition of kangaroo (128bit LSB) - int128_t d; // Travelled distance (b127=sign b126=kangaroo type, b125..b0 distance + int256_t d; // Travelled distance (b255=sign b254=kangaroo type, b253..b0 distance } ENTRY; +// On-disk and on-wire size of one ENTRY. Hard-coded rather than sizeof() at +// each use site so a layout change cannot silently reinterpret old files. +#define ENTRY_SIZE 48 +static_assert(sizeof(ENTRY) == ENTRY_SIZE,"ENTRY must stay packed at 48 bytes"); + typedef struct { uint32_t nbItem; @@ -69,7 +96,7 @@ class HashTable { HashTable(); int Add(Int *x,Int *d,uint32_t type); - int Add(uint64_t h,int128_t *x,int128_t *d); + int Add(uint64_t h,int128_t *x,int256_t *d); int Add(uint64_t h,ENTRY *e); uint64_t GetNbItem(); void Reset(); @@ -88,15 +115,16 @@ class HashTable { Int kDist; uint32_t kType; - static void Convert(Int *x,Int *d,uint32_t type,uint64_t *h,int128_t *X,int128_t *D); + static void Convert(Int *x,Int *d,uint32_t type,uint64_t *h,int128_t *X,int256_t *D); static int MergeH(uint32_t h,FILE* f1,FILE* f2,FILE* fd,uint32_t *nbDP,uint32_t* duplicate, Int* d1,uint32_t* k1,Int* d2,uint32_t* k2); - static void CalcDistAndType(int128_t d,Int* kDist,uint32_t* kType); + static void CalcDistAndType(int256_t d,Int* kDist,uint32_t* kType); private: - ENTRY *CreateEntry(int128_t *x,int128_t *d); + ENTRY *CreateEntry(int128_t *x,int256_t *d); static int compare(int128_t *i1,int128_t *i2); + static bool sameDist(int256_t *a,int256_t *b); std::string GetStr(int128_t *i); }; diff --git a/Kangaroo.cpp b/Kangaroo.cpp index 33104744..747c3ba9 100644 --- a/Kangaroo.cpp +++ b/Kangaroo.cpp @@ -313,7 +313,7 @@ bool Kangaroo::AddToTable(Int *pos,Int *dist,uint32_t kType) { } -bool Kangaroo::AddToTable(uint64_t h,int128_t *x,int128_t *d) { +bool Kangaroo::AddToTable(uint64_t h,int128_t *x,int256_t *d) { int addStatus = hashTable.Add(h,x,d); if(addStatus== ADD_COLLISION) { diff --git a/Kangaroo.h b/Kangaroo.h index 70f5e342..6d93b415 100644 --- a/Kangaroo.h +++ b/Kangaroo.h @@ -96,7 +96,7 @@ typedef struct { uint32_t kIdx; uint32_t h; int128_t x; - int128_t d; + int256_t d; } DP; @@ -117,9 +117,12 @@ typedef struct { } DP_CACHE; // Work file type -#define HEADW 0xFA6A8001 // Full work file -#define HEADK 0xFA6A8002 // Kangaroo only file -#define HEADKS 0xFA6A8003 // Compressed Kangaroo only file +// Bumped for the 256bit distance field: ENTRY grew from 32 to 48 bytes, so a +// pre-existing work file read with the new layout would silently desynchronise +// and produce garbage DPs. Old files are now rejected by magic instead. +#define HEADW 0xFA6A8011 // Full work file (256bit distance) +#define HEADK 0xFA6A8012 // Kangaroo only file (256bit distance) +#define HEADKS 0xFA6A8013 // Compressed Kangaroo only file (256bit distance) // Number of Hash entry per partition #define H_PER_PART (HASH_SIZE / MERGE_PART) @@ -165,7 +168,7 @@ class Kangaroo { void SetDP(int size); void CreateHerd(int nbKangaroo,Int *px, Int *py, Int *d, int firstType,bool lock=true); void CreateJumpTable(); - bool AddToTable(uint64_t h,int128_t *x,int128_t *d); + bool AddToTable(uint64_t h,int128_t *x,int256_t *d); bool AddToTable(Int *pos,Int *dist,uint32_t kType); bool SendToServer(std::vector &dp,uint32_t threadId,uint32_t gpuId); bool CheckKey(Int d1,Int d2,uint8_t type); @@ -181,7 +184,7 @@ class Kangaroo { void SaveWork(uint64_t totalCount,double totalTime,TH_PARAM *threads,int nbThread); void SaveServerWork(); void FetchWalks(uint64_t nbWalk,Int *x,Int *y,Int *d); - void FetchWalks(uint64_t nbWalk,std::vector& kangs,Int* x,Int* y,Int* d); + void FetchWalks(uint64_t nbWalk,std::vector& kangs,Int* x,Int* y,Int* d); void FectchKangaroos(TH_PARAM *threads); FILE *ReadHeader(std::string fileName,uint32_t *version,int type); bool SaveHeader(std::string fileName,FILE* f,int type,uint64_t totalCount,double totalTime); @@ -204,8 +207,8 @@ class Kangaroo { void InitSocket(); void WaitForServer(); int32_t GetServerStatus(); - bool SendKangaroosToServer(std::string& fileName,std::vector& kangs); - bool GetKangaroosFromServer(std::string& fileName,std::vector& kangs); + bool SendKangaroosToServer(std::string& fileName,std::vector& kangs); + bool GetKangaroosFromServer(std::string& fileName,std::vector& kangs); #ifdef WIN64 HANDLE ghMutex; diff --git a/Network.cpp b/Network.cpp index e6fc10ff..40b2d2f0 100644 --- a/Network.cpp +++ b/Network.cpp @@ -44,7 +44,9 @@ static SOCKET serverSock = 0; #define SERVER_VERSION 3 -#define SERVER_HEADER 0x67DEDDC1 +// Bumped with the 256bit distance: DP grew 40 -> 56 bytes and the kangaroo +// block grew 16 -> 32 bytes per kangaroo. An old peer would misparse both. +#define SERVER_HEADER 0x67DEDDD1 #define KANG_PER_BLOCK 2048 @@ -336,7 +338,7 @@ bool Kangaroo::HandleRequest(TH_PARAM *p) { uint64_t nbKangaroo = 0; uint32_t strSize; char fileName[256]; - int128_t* KBuff; + int256_t* KBuff; uint32_t nbK; uint32_t header = HEADKS; uint32_t version = 0; @@ -378,7 +380,7 @@ bool Kangaroo::HandleRequest(TH_PARAM *p) { PUT("nbKangaroo",p->clientSock,&nbKangaroo,sizeof(uint64_t),ntimeout); checkSum.SetInt32(0); - KBuff = (int128_t*)malloc(KANG_PER_BLOCK * sizeof(int128_t)); + KBuff = (int256_t*)malloc(KANG_PER_BLOCK * sizeof(int256_t)); while(nbKangaroo > 0) { @@ -389,15 +391,17 @@ bool Kangaroo::HandleRequest(TH_PARAM *p) { } for(uint32_t k = 0; k < nbK; k++) { - ::fread(&KBuff[k],16,1,f); + ::fread(&KBuff[k],32,1,f); // Checksum K.SetInt32(0); K.bits64[1] = KBuff[k].i64[1]; K.bits64[0] = KBuff[k].i64[0]; + K.bits64[2] = KBuff[k].i64[2]; + K.bits64[3] = KBuff[k].i64[3]; checkSum.Add(&K); } - PUTFREE("packet",p->clientSock,KBuff,nbK * 16,ntimeout,KBuff); + PUTFREE("packet",p->clientSock,KBuff,nbK * 32,ntimeout,KBuff); nbKangaroo -= nbK; @@ -422,7 +426,7 @@ bool Kangaroo::HandleRequest(TH_PARAM *p) { uint32_t fileNameSize; char fileNameTmp[264]; char fileName[256]; - int128_t *KBuff; + int256_t *KBuff; uint32_t nbK; uint32_t header = HEADKS; uint32_t version = 0; @@ -459,7 +463,7 @@ bool Kangaroo::HandleRequest(TH_PARAM *p) { checkSum.SetInt32(0); - KBuff = (int128_t *)malloc(KANG_PER_BLOCK*sizeof(int128_t)); + KBuff = (int256_t *)malloc(KANG_PER_BLOCK*sizeof(int256_t)); while(nbKangaroo>0) { @@ -469,14 +473,16 @@ bool Kangaroo::HandleRequest(TH_PARAM *p) { nbK = (uint32_t)nbKangaroo; } - GETFREE("packet",p->clientSock,KBuff,nbK * 16,ntimeout,KBuff); + GETFREE("packet",p->clientSock,KBuff,nbK * 32,ntimeout,KBuff); for(uint32_t k = 0; k < nbK; k++) { - ::fwrite(&KBuff[k],16,1,f); + ::fwrite(&KBuff[k],32,1,f); // Checksum K.SetInt32(0); K.bits64[1] = KBuff[k].i64[1]; K.bits64[0] = KBuff[k].i64[0]; + K.bits64[2] = KBuff[k].i64[2]; + K.bits64[3] = KBuff[k].i64[3]; checkSum.Add(&K); } @@ -696,7 +702,7 @@ void Kangaroo::RunServer() { } SetDP(initDPSize); - if(sizeof(DP) != 40) { + if(sizeof(DP) != 56) { ::printf("Error: Invalid DP size struct\n"); exit(-1); } @@ -980,14 +986,14 @@ void Kangaroo::WaitForServer() { } // Get Kangaroo from server -bool Kangaroo::GetKangaroosFromServer(std::string& fileName,std::vector& kangs) { +bool Kangaroo::GetKangaroosFromServer(std::string& fileName,std::vector& kangs) { int nbRead; int nbWrite; uint32_t fileNameSize = (uint32_t)fileName.length(); uint64_t nbKangaroo = 0; uint32_t nbK; - int128_t* KBuff; + int256_t* KBuff; Int checkSum; WaitForServer(); @@ -1011,7 +1017,7 @@ bool Kangaroo::GetKangaroosFromServer(std::string& fileName,std::vector& kangs) { +bool Kangaroo::SendKangaroosToServer(std::string& fileName,std::vector& kangs) { int nbWrite; uint32_t fileNameSize = (uint32_t)fileName.length(); uint64_t nbKangaroo = kangs.size(); uint64_t pos; uint32_t nbK; - int128_t *KBuff; + int256_t *KBuff; Int checkSum; WaitForServer(); @@ -1088,7 +1096,7 @@ bool Kangaroo::SendKangaroosToServer(std::string& fileName,std::vector PUT("fileName",serverConn,fileName.c_str(),fileNameSize,ntimeout); PUT("nbKangaroo",serverConn,&nbKangaroo,sizeof(uint64_t),ntimeout); - KBuff = (int128_t*)malloc(KANG_PER_BLOCK * sizeof(int128_t)); + KBuff = (int256_t*)malloc(KANG_PER_BLOCK * sizeof(int256_t)); checkSum.SetInt32(0); pos = 0; @@ -1107,17 +1115,19 @@ bool Kangaroo::SendKangaroosToServer(std::string& fileName,std::vector } for(uint32_t k = 0; k < nbK; k++) { - memcpy(&KBuff[k],&kangs[pos],16); + memcpy(&KBuff[k],&kangs[pos],32); pos++; // Checksum Int K; K.SetInt32(0); K.bits64[1] = KBuff[k].i64[1]; K.bits64[0] = KBuff[k].i64[0]; + K.bits64[2] = KBuff[k].i64[2]; + K.bits64[3] = KBuff[k].i64[3]; checkSum.Add(&K); } - PUTFREE("packet",serverConn,KBuff,nbK * 16,ntimeout,KBuff); + PUTFREE("packet",serverConn,KBuff,nbK * 32,ntimeout,KBuff); nbKangaroo -= nbK; @@ -1154,7 +1164,7 @@ bool Kangaroo::SendToServer(std::vector &dps,uint32_t threadId,uint32_t gp for(uint32_t i = 0; i &dps,uint32_t threadId,uint32_t gp dp[i].x.i64[1] = X.i64[1]; dp[i].d.i64[0] = D.i64[0]; dp[i].d.i64[1] = D.i64[1]; + dp[i].d.i64[2] = D.i64[2]; + dp[i].d.i64[3] = D.i64[3]; } diff --git a/build_cpu.bat b/build_cpu.bat new file mode 100644 index 00000000..0d8fc1ea --- /dev/null +++ b/build_cpu.bat @@ -0,0 +1,14 @@ +@echo off +setlocal +call "C:\Program Files (x86)\Microsoft Visual Studio\18\BuildTools\VC\Auxiliary\Build\vcvars64.bat" >nul +if errorlevel 1 (echo VCVARS FAILED & exit /b 1) +if not exist obj mkdir obj +cl /nologo /EHsc /O2 /std:c++14 /DWIN64 /D_CRT_SECURE_NO_WARNINGS /I. /c ^ + main.cpp Kangaroo.cpp HashTable.cpp Backup.cpp Thread.cpp Check.cpp ^ + Network.cpp Merge.cpp PartMerge.cpp Timer.cpp ^ + SECPK1\Int.cpp SECPK1\IntGroup.cpp SECPK1\IntMod.cpp ^ + SECPK1\Point.cpp SECPK1\Random.cpp SECPK1\SECP256K1.cpp /Foobj\ +if errorlevel 1 (echo COMPILE FAILED & exit /b 1) +link /nologo /OUT:kangaroo.exe obj\*.obj ws2_32.lib advapi32.lib +if errorlevel 1 (echo LINK FAILED & exit /b 1) +echo BUILD OK diff --git a/build_gpu.bat b/build_gpu.bat new file mode 100644 index 00000000..eeed336d --- /dev/null +++ b/build_gpu.bat @@ -0,0 +1,47 @@ +@echo off +setlocal enabledelayedexpansion +rem GPU build: nvcc for GPUEngine.cu, MSVC for the rest, -DWITHGPU throughout. +rem sm_86 = GA107 (RTX 3050 Laptop). Override with: build_gpu.bat 89 + +set CCAP=%1 +if "%CCAP%"=="" set CCAP=86 + +rem CUDA 13.x may reject the newest MSVC as an unsupported host compiler. +rem MSVC_DIR pins nvcc to a known-good toolset; leave empty to use the default. +set "MSVC_DIR=C:\Program Files (x86)\Microsoft Visual Studio\18\BuildTools\VC\Tools\MSVC\14.44.35207\bin\Hostx64\x64" + +call "C:\Program Files (x86)\Microsoft Visual Studio\18\BuildTools\VC\Auxiliary\Build\vcvars64.bat" >nul +if errorlevel 1 (echo VCVARS FAILED & exit /b 1) + +for /f "delims=" %%i in ('dir /b /o-n "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v*" 2^>nul') do ( + if not defined CUDA_VER set CUDA_VER=%%i +) +if not defined CUDA_VER (echo CUDA TOOLKIT NOT FOUND & exit /b 1) +set "CUDA_PATH=C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\%CUDA_VER%" +echo Using %CUDA_VER%, sm_%CCAP% + +rem separate obj dir: the CPU build left non-WITHGPU objects in obj\ +if not exist objgpu mkdir objgpu +del /q objgpu\*.obj 2>nul + +set CCBIN= +if exist "%MSVC_DIR%\cl.exe" set CCBIN=-ccbin "%MSVC_DIR%" + +"%CUDA_PATH%\bin\nvcc.exe" -maxrregcount=0 --ptxas-options=-v --compile %CCBIN% ^ + -m64 -O2 -I. -I"%CUDA_PATH%\include" -DWITHGPU -DWIN64 -D_CRT_SECURE_NO_WARNINGS ^ + -gencode=arch=compute_%CCAP%,code=sm_%CCAP% ^ + -o objgpu\GPUEngine.obj -c GPU\GPUEngine.cu +if errorlevel 1 (echo NVCC FAILED & exit /b 1) + +cl /nologo /EHsc /O2 /std:c++14 /DWITHGPU /DWIN64 /D_CRT_SECURE_NO_WARNINGS ^ + /I. /I"%CUDA_PATH%\include" /c ^ + main.cpp Kangaroo.cpp HashTable.cpp Backup.cpp Thread.cpp Check.cpp ^ + Network.cpp Merge.cpp PartMerge.cpp Timer.cpp ^ + SECPK1\Int.cpp SECPK1\IntGroup.cpp SECPK1\IntMod.cpp ^ + SECPK1\Point.cpp SECPK1\Random.cpp SECPK1\SECP256K1.cpp /Foobjgpu\ +if errorlevel 1 (echo COMPILE FAILED & exit /b 1) + +link /nologo /OUT:kangaroo-gpu.exe objgpu\*.obj ws2_32.lib advapi32.lib ^ + /LIBPATH:"%CUDA_PATH%\lib\x64" cudart.lib +if errorlevel 1 (echo LINK FAILED & exit /b 1) +echo BUILD OK diff --git a/test140.txt b/test140.txt new file mode 100644 index 00000000..d47d9757 --- /dev/null +++ b/test140.txt @@ -0,0 +1,3 @@ +80000000000000000000000000000000000 +FFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFF +031F6A332D3C5C4F2DE2378C012F429CD109BA07D69690C6C701B6BB87860D6640 From e4a2d8ecccd51b9901341459953efa01feb20489 Mon Sep 17 00:00:00 2001 From: Dileep-Kumar-5 Date: Wed, 5 Aug 2026 23:12:01 +0530 Subject: [PATCH 2/7] fix: build against CUDA 13 and locate a non-admin toolkit Two changes needed to actually compile the GPU path. cudaDeviceProp::computeMode was removed in CUDA 13, so GPUEngine.cu no longer compiled at all -- this is upstream breakage, unrelated to the 256-bit distance work. The cudaDevAttrComputeMode device attribute still exists, so query that instead and keep the same startup line. Guard the index since the attribute is not bounded by the sComputeMode table. build_gpu.bat now prefers a toolkit extracted beside the repo before falling back to the Program Files install. Installing CUDA properly needs admin; NVIDIA's per-component redist archives (cuda_nvcc, cuda_cudart, cuda_crt, libnvvm -- 89 MB total vs ~9 GB) extract to a user-writable directory and are enough to build. Verified on an RTX 3050 Laptop (sm_86, CUDA 13.3.73, MSVC 14.44): - solves the 56-bit sample key correctly at 322 MK/s - on the 2^139 puzzle #140 range, 713,781 GPU-produced DPs at 100.000% -wcheck - ptxas: 98 registers, 0 spills, stack frame 18560 -> 20608 bytes - throughput 382.7 -> 366.7 MK/s, about 4% slower (single run each) Co-Authored-By: Claude Opus 5 --- .gitignore | 3 +++ GPU/GPUEngine.cu | 9 ++++++++- build_gpu.bat | 16 +++++++++++----- 3 files changed, 22 insertions(+), 6 deletions(-) diff --git a/.gitignore b/.gitignore index a89636b9..86e72863 100644 --- a/.gitignore +++ b/.gitignore @@ -21,3 +21,6 @@ objgpu/ # run artifacts *.kcp *.log + +# CUDA runtime staged next to the binary by build_gpu.bat +cudart64_*.dll diff --git a/GPU/GPUEngine.cu b/GPU/GPUEngine.cu index bfdc49e8..cb6fcba1 100644 --- a/GPU/GPUEngine.cu +++ b/GPU/GPUEngine.cu @@ -364,11 +364,18 @@ void GPUEngine::PrintCudaInfo() { cudaDeviceProp deviceProp; cudaGetDeviceProperties(&deviceProp,i); + + // cudaDeviceProp::computeMode was removed in CUDA 13; the device + // attribute still exists, so query that instead of the struct field. + int computeMode = 0; + cudaDeviceGetAttribute(&computeMode,cudaDevAttrComputeMode,i); + if(computeMode < 0 || computeMode > 3) computeMode = 4; // "Unknown" + printf("GPU #%d %s (%dx%d cores) (Cap %d.%d) (%.1f MB) (%s)\n", i,deviceProp.name,deviceProp.multiProcessorCount, _ConvertSMVer2Cores(deviceProp.major,deviceProp.minor), deviceProp.major,deviceProp.minor,(double)deviceProp.totalGlobalMem / 1048576.0, - sComputeMode[deviceProp.computeMode]); + sComputeMode[computeMode]); } diff --git a/build_gpu.bat b/build_gpu.bat index eeed336d..210d7d12 100644 --- a/build_gpu.bat +++ b/build_gpu.bat @@ -13,12 +13,18 @@ set "MSVC_DIR=C:\Program Files (x86)\Microsoft Visual Studio\18\BuildTools\VC\To call "C:\Program Files (x86)\Microsoft Visual Studio\18\BuildTools\VC\Auxiliary\Build\vcvars64.bat" >nul if errorlevel 1 (echo VCVARS FAILED & exit /b 1) -for /f "delims=" %%i in ('dir /b /o-n "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v*" 2^>nul') do ( - if not defined CUDA_VER set CUDA_VER=%%i +rem Prefer a local extracted toolkit (redist zips, no admin install), else +rem the system install under Program Files. +set "CUDA_PATH=%~dp0..\cuda-toolkit" +if not exist "%CUDA_PATH%\bin\nvcc.exe" ( + for /f "delims=" %%i in ('dir /b /o-n "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v*" 2^>nul') do ( + if not defined CUDA_VER set CUDA_VER=%%i + ) + if not defined CUDA_VER (echo CUDA TOOLKIT NOT FOUND & exit /b 1) + set "CUDA_PATH=C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\!CUDA_VER!" ) -if not defined CUDA_VER (echo CUDA TOOLKIT NOT FOUND & exit /b 1) -set "CUDA_PATH=C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\%CUDA_VER%" -echo Using %CUDA_VER%, sm_%CCAP% +if not exist "%CUDA_PATH%\bin\nvcc.exe" (echo nvcc NOT FOUND under %CUDA_PATH% & exit /b 1) +echo Using "%CUDA_PATH%", sm_%CCAP% rem separate obj dir: the CPU build left non-WITHGPU objects in obj\ if not exist objgpu mkdir objgpu From 10cde3f8142d0d5d11aee23ea869d7bbdcd3f61c Mon Sep 17 00:00:00 2001 From: Dileep-Kumar-5 Date: Thu, 6 Aug 2026 01:20:43 +0530 Subject: [PATCH 3/7] fix: truncate the kangaroo-transfer checksum to 256 bits The kangaroo backup never completed: the server writes .tmp and only renames it into place when the client's checksum matches, so every transfer silently produced a stale .tmp and no backup. checkSum is an Int (5 x 64 bits) but only its low 32 bytes are put on the wire, and the receiver compares with IsEqual, which tests all five words. With the old 126-bit distances a herd summed to well under 2^256 and bits64[4] stayed zero, so nobody noticed. The 256-bit distance puts the kangaroo type at b254 and the sign at b255, so ~2^19 kangaroos overflow past 2^256 and bits64[4] no longer matches the freshly-zeroed receiver. Define the checksum as mod 2^256 -- what the wire actually carries -- and clear bits64[4] on both sides before sending and before comparing. All four sites: server send/compare and client send/compare. check_layout.py now asserts all four sites truncate before use. Verified on localhost, 2^139 range, GPU client: - DP wire: 196,231 DPs arrive over TCP, 100.000% -wcheck - upload: kang written and renamed, 17,170,448 bytes (536,576 x 32) - download: "2^19.03 kangaroos loaded, 0 created" - DPs generated from downloaded kangaroos: 216,445, 100.000% -wcheck Co-Authored-By: Claude Opus 5 --- Network.cpp | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/Network.cpp b/Network.cpp index 40b2d2f0..baf34bce 100644 --- a/Network.cpp +++ b/Network.cpp @@ -409,6 +409,9 @@ bool Kangaroo::HandleRequest(TH_PARAM *p) { free(KBuff); + // checksum is defined mod 2^256: only 32 bytes go on the wire, and the + // distance flags at b254/b255 make a herd overflow past 256 bits. + checkSum.bits64[4] = 0; PUT("checkSum",p->clientSock,checkSum.bits64,32,ntimeout); ::fclose(f); @@ -496,6 +499,9 @@ bool Kangaroo::HandleRequest(TH_PARAM *p) { K.SetInt32(0); GET("checksum",p->clientSock,K.bits64,32,ntimeout); + // checksum is defined mod 2^256: only 32 bytes go on the wire, and the + // distance flags at b254/b255 make a herd overflow past 256 bits. + checkSum.bits64[4] = 0; if(!K.IsEqual(&checkSum)) { ::printf("\nWarning, Kangaroo backup wrong checksum %s\n",fileName); } else { @@ -1059,6 +1065,9 @@ bool Kangaroo::GetKangaroosFromServer(std::string& fileName,std::vector free(KBuff); + // checksum is defined mod 2^256: only 32 bytes go on the wire, and the + // distance flags at b254/b255 make a herd overflow past 256 bits. + checkSum.bits64[4] = 0; PUT("checksum",serverConn,checkSum.bits64,32,ntimeout); } From 6acfc06d1ef542fd7be4b1209c99a36af57883d0 Mon Sep 17 00:00:00 2001 From: Dileep-Kumar-5 Date: Thu, 6 Aug 2026 01:21:16 +0530 Subject: [PATCH 4/7] test: add the layout and distance-field verification scripts These are what the previous two commits refer to; they lived outside the repo until now. tools/check_layout.py parses the sources and asserts producer/consumer agreement on everything the widening touched by hand: each ITEM uint32 written exactly once and read at the same offset, device slot usage bounded by KSIZE, the Add256 carry chain order, DP struct size against its own runtime assertion, all four checksum sites truncating before use, and both format magics bumped. It caught two real defects during this work -- KBuff mallocs still cast to int128_t*, and an over-broad rule of its own. tools/test_distfield.py models Convert/CalcDistAndType mask-for-mask and checks sign/type/magnitude round-trip at 139..253 bits, that a 254-bit magnitude is rejected rather than truncated, and that the flag bits are disjoint from the magnitude. Neither needs a compiler or a GPU: python tools/check_layout.py python tools/test_distfield.py Co-Authored-By: Claude Opus 5 --- .gitignore | 5 + tools/check_layout.py | 225 ++++++++++++++++++++++++++++++++++++++++ tools/test_distfield.py | 116 +++++++++++++++++++++ 3 files changed, 346 insertions(+) create mode 100644 tools/check_layout.py create mode 100644 tools/test_distfield.py diff --git a/.gitignore b/.gitignore index 86e72863..765b6072 100644 --- a/.gitignore +++ b/.gitignore @@ -24,3 +24,8 @@ objgpu/ # CUDA runtime staged next to the binary by build_gpu.bat cudart64_*.dll + +# kangaroo backup files written by -wss +kang +kang.tmp +*.tmp diff --git a/tools/check_layout.py b/tools/check_layout.py new file mode 100644 index 00000000..4f78e151 --- /dev/null +++ b/tools/check_layout.py @@ -0,0 +1,225 @@ +#!/usr/bin/env python3 +"""Static consistency check for the 256-bit distance patch. + +Phase 2 is mostly hand-edited offsets across the CUDA kernel, the host readout +and two wire formats. A compiler catches type errors; it does NOT catch a DP +written at word 9 and read at word 8. This parses the actual sources and +asserts producer and consumer agree. +""" + +import re +import sys +from pathlib import Path + +# Works both from inside the repo (tools/) and from a parent dir holding +# a "Kangaroo" checkout. +_here = Path(__file__).resolve().parent +ROOT = _here.parent if (_here.parent / "HashTable.h").exists() else _here.parent / "Kangaroo" + + +def read(rel): + return (ROOT / rel).read_text(encoding="utf-8", errors="replace") + + +def define(src, name): + m = re.search(r"^#define\s+%s\s+(\d+)" % re.escape(name), src, re.M) + assert m, "no #define %s" % name + return int(m.group(1)) + + +def check_item_layout(): + """GPU DP output: OutputDP writer vs GPUEngine.cu reader.""" + eng_h = read("GPU/GPUEngine.h") + item_size = define(eng_h, "ITEM_SIZE") + item32 = item_size // 4 + + # x[4 u64] + d[4 u64] + kIdx[1 u64] + assert item_size == (4 + 4 + 1) * 8 == 72, item_size + + math_h = read("GPU/GPUMath.h") + body = math_h[math_h.index("#define OutputDP"):] + body = body[:body.index("\n}")] + slots = [(int(o), field, int(i)) + for o, field, i in re.findall( + r"out\[pos\*ITEM_SIZE32 \+ (\d+)\] = \(\(uint32_t \*\)(\w+)\)\[(\d+)\];", + body)] + assert slots, "OutputDP body not parsed" + + # Every uint32 of the item is written exactly once, contiguous from 1. + written = [s[0] for s in slots] + assert written == list(range(1, item32 + 1)), \ + "OutputDP writes %s, expected 1..%d" % (written, item32) + + # Each source field is written low word first, no gaps. + for field, count in (("x", 8), ("d", 8), ("idx", 2)): + idxs = [i for _, f, i in slots if f == field] + assert idxs == list(range(count)), "%s writes %s" % (field, idxs) + + # Reader side: itemPtr = out + i*ITEM_SIZE32 + 1, so writer slot N -> itemPtr[N-1]. + eng_cu = read("GPU/GPUEngine.cu") + x_off = [int(n) for n in re.findall(r"it\.x\.bits64\[(\d)\] = x\[\d\];", eng_cu)] + d_base = int(re.search(r"uint64_t \*d = \(uint64_t \*\)\(itemPtr \+ (\d+)\);", + eng_cu).group(1)) + d_words = re.findall(r"it\.d\.bits64\[(\d)\] = d\[(\d)\];", eng_cu) + k_off = int(re.search(r"it\.kIdx = \*\(\(uint64_t\*\)\(itemPtr \+ (\d+)\)\);", + eng_cu).group(1)) + + assert x_off == [0, 1, 2, 3], x_off + assert d_base == 8, "d read at u32 offset %d, writer put it at 8" % d_base + assert [(int(a), int(b)) for a, b in d_words] == [(0, 0), (1, 1), (2, 2), (3, 3)], \ + "distance words not read 1:1: %s" % d_words + assert k_off == 16, "kIdx read at u32 %d, writer put it at 16" % k_off + # kIdx is the last 2 u32 of the item + assert k_off + 2 == item32, "kIdx at %d + 2 != item32 %d" % (k_off, item32) + return item_size + + +def check_kangaroo_slots(): + """Device kangaroo record: every slot used must fit inside KSIZE.""" + eng_h = read("GPU/GPUEngine.h") + ksize_sym = int(re.search(r"#ifdef USE_SYMMETRY\s*\n#define KSIZE (\d+)", eng_h).group(1)) + ksize = int(re.search(r"#else\s*\n#define KSIZE (\d+)", eng_h).group(1)) + assert (ksize, ksize_sym) == (12, 13), (ksize, ksize_sym) + + math_h = read("GPU/GPUMath.h") + used = {int(n) for n in re.findall(r"IDX \+ (\d+) \* blockDim\.x \+ stride", math_h)} + assert used == set(range(13)), "device slots used: %s" % sorted(used) + assert max(used) < ksize_sym + + # Non-symmetry build must not touch slot 12 (that is the lastJump word). + non_sym = re.sub(r"#ifdef USE_SYMMETRY.*?#endif", "", math_h, flags=re.S) + used_ns = {int(n) for n in re.findall(r"IDX \+ (\d+) \* blockDim\.x \+ stride", non_sym)} + assert max(used_ns) < ksize, "non-symmetry uses slot %d, KSIZE=%d" % (max(used_ns), ksize) + + # Host side writes the same slots. + eng_cu = read("GPU/GPUEngine.cu") + host = {int(n) for n in re.findall(r"t \+ (\d+) \* nbThreadPerGroup", eng_cu)} + assert host == set(range(13)), "host slots: %s" % sorted(host) + + # dist is 4 words everywhere it is declared. + assert "dist[GPU_GRP_SIZE][2]" not in math_h + assert "dist[GPU_GRP_SIZE][2]" not in read("GPU/GPUCompute.h") + # and accumulated with the 4-word carry chain + assert "Add256(dist[g],jD[jmp]);" in read("GPU/GPUCompute.h") + assert "jD[NB_JUMP][4]" in math_h + return ksize + + +def check_add256(): + """Carry must chain through all four words: add.cc / addc.cc / addc.cc / addc.""" + math_h = read("GPU/GPUMath.h") + body = math_h[math_h.index("#define Add256"):] + body = body[:body.index("(a)[3]);}") + len("(a)[3]);}")] + ops = re.findall(r"(UADDO1|UADDC1|UADD1)\(\(r\)\[(\d)\], \(a\)\[(\d)\]\)", body) + assert [o[0] for o in ops] == ["UADDO1", "UADDC1", "UADDC1", "UADD1"], ops + assert [o[1] for o in ops] == list("0123") and [o[2] for o in ops] == list("0123") + + +def check_wire(): + """DP packet and kangaroo block sizes agree with their runtime assertions.""" + kh = read("Kangaroo.h") + dp = kh[kh.index("// DP transfered over the network"):] + dp = dp[:dp.index("} DP;")] + assert "int128_t x;" in dp and "int256_t d;" in dp, dp + dp_size = 4 + 4 + 16 + 32 # kIdx + h + x + d + assert dp_size == 56 + + net = read("Network.cpp") + asserted = int(re.search(r"if\(sizeof\(DP\) != (\d+)\)", net).group(1)) + assert asserted == dp_size, "runtime check says %d, struct is %d" % (asserted, dp_size) + + # All four distance words are packed into the outgoing DP. + packed = re.findall(r"dp\[i\]\.d\.i64\[(\d)\] = D\.i64\[(\d)\];", net) + assert [(int(a), int(b)) for a, b in packed] == [(0, 0), (1, 1), (2, 2), (3, 3)], packed + + # Kangaroo block: 32 bytes per kangaroo on every path, none left at 16. + assert "int128_t* KBuff" not in net and "int128_t *KBuff" not in net + # Declaration, cast and allocation size must all agree -- a stale cast here + # is a type error the eye slides right over. + allocs = re.findall(r"KBuff = \((\w+) ?\*\)malloc\(KANG_PER_BLOCK ?\* ?sizeof\((\w+)\)\)", net) + assert len(allocs) == 4, allocs + assert all(c == "int256_t" and t == "int256_t" for c, t in allocs), allocs + for pat in (r"::fread\(&KBuff\[k\],(\d+),1,f\);", + r"::fwrite\(&KBuff\[k\],(\d+),1,f\);", + r"memcpy\(&KBuff\[k\],&kangs\[pos\],(\d+)\);"): + sizes = re.findall(pat, net) + assert sizes and all(s == "32" for s in sizes), (pat, sizes) + pkt = re.findall(r"KBuff,nbK \* (\d+),ntimeout", net) + assert pkt and all(s == "32" for s in pkt), pkt + + # Checksum must cover all four words wherever it is computed. + blocks = re.findall(r"K\.SetInt32\(0\);(.*?)checkSum\.Add\(&K\);", net, re.S) + assert blocks, "no checksum blocks found" + for b in blocks: + got = sorted(int(n) for n in re.findall(r"K\.bits64\[(\d)\] = KBuff", b)) + assert got == [0, 1, 2, 3], "checksum covers words %s" % got + + # The checksum accumulator is a 5-word Int but only 32 bytes go on the wire, + # and the distance now carries flags at b254/b255, so a herd overflows past + # 256 bits. Both sides must truncate or every transfer fails its checksum. + sends = len(re.findall(r'PUT\("check[Ss]um",\w+(?:->clientSock)?,checkSum\.bits64,32', net)) + compares = len(re.findall(r"if\(!K\.IsEqual\(&checkSum\)\)", net)) + truncs = len(re.findall(r"checkSum\.bits64\[4\] = 0;", net)) + assert sends + compares == 4, (sends, compares) + assert truncs == 4, "checksum truncated at %d of 4 sites" % truncs + for m in re.finditer(r"checkSum\.bits64\[4\] = 0;(.{0,400})", net, re.S): + assert ("IsEqual(&checkSum)" in m.group(1) + or "checkSum.bits64,32" in m.group(1)), "truncation not before use" + + assert "kangs.size()*32 + 16" in read("Backup.cpp") + return dp_size, len(blocks) + + +def check_entry(): + ht = read("HashTable.h") + assert define(ht, "ENTRY_SIZE") == 16 + 32 == 48 + assert define(ht, "MAX_INTERVAL_BITS") == 253 + assert "int256_t d;" in ht + # Every ENTRY-sized transfer goes through ENTRY_SIZE, and the two halves of + # an ENTRY are read/written at their own widths. Bare 32s elsewhere in these + # files are 256-bit Int header fields (range start/end, key x/y) -- correct + # as-is, so this checks the ENTRY sites specifically rather than banning 32. + htc = read("HashTable.cpp") + assert "::fread(items+i,ENTRY_SIZE,1,f);" in read("Check.cpp") + for pat in (r"::fread\(&e1,ENTRY_SIZE,1,f1\)", r"::fread\(&e2,ENTRY_SIZE,1,f2\)", + r"::fwrite\(output,ENTRY_SIZE,nbd,fd\)", + r"uint64_t hSize = \(uint64_t\)ENTRY_SIZE \* E\[h\]\.nbItem;"): + assert re.search(pat, htc), pat + assert len(re.findall(r"memcpy\(output ?\+ ?nbd,&e\d,ENTRY_SIZE\)", htc)) == 5 + for half, width in (("x", "int128_t"), ("d", "int256_t")): + assert ("fwrite(&(E[h].items[i]->%s),sizeof(%s),1,f)" % (half, width)) in htc + assert ("fread(&(e->%s),sizeof(%s),1,f)" % (half, width)) in htc + # The truncating mask is gone; the guard replaced it. + htc = read("HashTable.cpp") + assert "exit(-1);" in htc and "0xC000000000000000ULL" in htc + + +def check_format_magic(): + """Old files and old peers must be rejected, not misparsed.""" + kh = read("Kangaroo.h") + for name, old in (("HEADW", "0xFA6A8001"), ("HEADK", "0xFA6A8002"), + ("HEADKS", "0xFA6A8003")): + m = re.search(r"#define %s\s+(0x[0-9A-Fa-f]+)" % name, kh) + assert m and m.group(1).lower() != old.lower(), "%s not bumped" % name + net = read("Network.cpp") + m = re.search(r"#define SERVER_HEADER (0x[0-9A-Fa-f]+)", net) + assert m and m.group(1).lower() != "0x67deddc1", "SERVER_HEADER not bumped" + + +def main(): + item = check_item_layout() + ksize = check_kangaroo_slots() + check_add256() + dp_size, nblocks = check_wire() + check_entry() + check_format_magic() + print("layout self-check OK") + print(" ENTRY 48 bytes (was 32) x:16 d:32") + print(" ITEM %d bytes (was 56) x:32 d:32 kIdx:8" % item) + print(" KSIZE %d words (was 10) px:4 py:4 dist:4" % ksize) + print(" DP packet %d bytes (was 40)" % dp_size) + print(" kangaroo 32 bytes (was 16), %d checksum sites widened" % nblocks) + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/tools/test_distfield.py b/tools/test_distfield.py new file mode 100644 index 00000000..fa150b1d --- /dev/null +++ b/tools/test_distfield.py @@ -0,0 +1,116 @@ +#!/usr/bin/env python3 +"""Model of the patched HashTable distance packing, to check the bit masks. + +Mirrors HashTable::Convert / CalcDistAndType word-for-word so the masks can be +exercised without a C++ toolchain. If this fails, the C++ is wrong too. +""" + +N_ORDER = 0xFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFEBAAEDCE6AF48A03BBFD25E8CD0364141 + +SIGN = 0x8000000000000000 # b255 of the 256-bit field, i.e. b63 of word 3 +TYPE = 0x4000000000000000 # b254 +MAG3 = 0x3FFFFFFFFFFFFFFF # b253..b192 + +MAX_INTERVAL_BITS = 253 + + +class Overflow(Exception): + pass + + +def words(v, n): + return [(v >> (64 * i)) & 0xFFFFFFFFFFFFFFFF for i in range(n)] + + +def convert256(d_signed, type_bit): + """HashTable::Convert, 256-bit path. d_signed is the true signed distance.""" + d = d_signed % N_ORDER # how Int holds it: mod n + w = words(d, 4) + sign = 0 + if w[3] > 0x7FFFFFFFFFFFFFFF: # upper half means negative + d = (-d_signed) % N_ORDER # ModNegK1order + w = words(d, 4) + sign = SIGN + if w[3] & 0xC000000000000000: + raise Overflow("distance exceeds 254 bits") + return [w[0], w[1], w[2], (w[3] & MAG3) | sign | (type_bit << 62)] + + +def calc_dist_and_type256(D): + """HashTable::CalcDistAndType, 256-bit path.""" + ktype = 1 if (D[3] & TYPE) else 0 + sign = 1 if (D[3] & SIGN) else 0 + mag = D[0] | (D[1] << 64) | (D[2] << 128) | ((D[3] & MAG3) << 192) + return ((-mag) % N_ORDER if sign else mag), ktype + + +def convert128(d_signed, type_bit): + """Legacy 126-bit path, now guarded instead of truncating.""" + d = d_signed % N_ORDER + w = words(d, 4) + sign = 0 + if w[3] > 0x7FFFFFFFFFFFFFFF: + d = (-d_signed) % N_ORDER + w = words(d, 4) + sign = 1 << 63 + if w[3] or w[2] or (w[1] & 0xC000000000000000): + raise Overflow("distance exceeds 126 bits") + return [w[0], (w[1] & 0x3FFFFFFFFFFFFFFF) | sign | (type_bit << 62)] + + +def widen(D128): + """HashTable::Widen -- legacy 126-bit entry into the 254-bit field.""" + return [D128[0], D128[1] & 0x3FFFFFFFFFFFFFFF, 0, + D128[1] & 0xC000000000000000] + + +def roundtrip(d_signed, type_bit): + got, ktype = calc_dist_and_type256(convert256(d_signed, type_bit)) + return got == d_signed % N_ORDER and ktype == type_bit + + +def demo(): + # Puzzle #140 lives in [2^139, 2^140); walks overshoot, so exercise well past it. + for bits in (139, 140, 160, 200, 253): + for t in (0, 1): + v = (1 << bits) - 12345 + assert roundtrip(v, t), (bits, t, "positive") + assert roundtrip(-v, t), (bits, t, "negative") + + # Boundary: 253 bits of magnitude fits, 254 does not. + assert roundtrip((1 << 253) - 1, 0) + try: + convert256(1 << 254, 0) + raise AssertionError("254-bit magnitude must be rejected, not truncated") + except Overflow: + pass + + # The original bug: a #140-scale distance silently lost its high bits in the + # 126-bit field. Now it raises instead of returning a wrong key. + try: + convert128(1 << 139, 0) + raise AssertionError("126-bit path must reject a 139-bit distance") + except Overflow: + pass + + # Legacy entries still decode correctly once widened. + for v in (1, 42, (1 << 125) - 1): + for t in (0, 1): + got, ktype = calc_dist_and_type256(widen(convert128(v, t))) + assert got == v and ktype == t, (v, t, got, ktype) + got, ktype = calc_dist_and_type256(widen(convert128(-v, t))) + assert got == (-v) % N_ORDER and ktype == t, (v, t) + + # Flags must never collide with magnitude bits. + assert SIGN | TYPE | MAG3 == 0xFFFFFFFFFFFFFFFF + assert SIGN & MAG3 == 0 and TYPE & MAG3 == 0 and SIGN & TYPE == 0 + + # ENTRY layout: 16-byte x + 32-byte d, no padding. + assert 16 + 32 == 48 + + print("distance-field self-check OK (254-bit magnitude, max interval %d bits)" + % MAX_INTERVAL_BITS) + + +if __name__ == "__main__": + demo() From 3a95e0b78431a5971b17dd399bc8c9d4fbcb6bed Mon Sep 17 00:00:00 2001 From: Dileep-Kumar-5 Date: Thu, 6 Aug 2026 01:37:15 +0530 Subject: [PATCH 5/7] test: add target verification and puzzle 120/135 inputs tools/verify_target.py derives the P2PKH address from a target's compressed pubkey and checks it against the published puzzle address, plus that the range is the canonical [2^(n-1), 2^n-1]. A pubkey pasted from a forum that does not belong to the address makes an entire run worthless and nothing in the solver would ever say so. Self-check derives puzzle #105's address from its published private key. tools/kangaroo_est.py estimates runtime, cost and DP storage. It now flags puzzles that are already solved -- #120 prints a 2.67x "profit" that is pure fiction because the coins were swept years ago -- and rejects a custom GPU profile that is missing --watts instead of crashing on None. Verified both builds against puzzle #120 (2^119, inside the old cap): stock 17,088 DPs 100.000% OK patched 15,795 DPs 100.000% OK versus puzzle #140 (2^139, outside it): stock 0.024% OK patched 100.000% OK So the widening fixes the >125-bit case without disturbing the case stock already handled. Co-Authored-By: Claude Opus 5 --- test120.txt | 3 + test135.txt | 3 + tools/kangaroo_est.py | 284 +++++++++++++++++++++++++++++++++++++++++ tools/verify_target.py | 152 ++++++++++++++++++++++ 4 files changed, 442 insertions(+) create mode 100644 test120.txt create mode 100644 test135.txt create mode 100644 tools/kangaroo_est.py create mode 100644 tools/verify_target.py diff --git a/test120.txt b/test120.txt new file mode 100644 index 00000000..5af10a2f --- /dev/null +++ b/test120.txt @@ -0,0 +1,3 @@ +800000000000000000000000000000 +FFFFFFFFFFFFFFFFFFFFFFFFFFFFFF +02CEB6CBBCDBDF5EF7150682150F4CE2C6F4807B349827DCDBDD1F2EFA885A2630 diff --git a/test135.txt b/test135.txt new file mode 100644 index 00000000..9a5e4cc6 --- /dev/null +++ b/test135.txt @@ -0,0 +1,3 @@ +4000000000000000000000000000000000 +7FFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFF +02145D2611C823A396EF6712CE0F712F09B9B4F3135E3E0AA3230FB9B6D08D1E16 diff --git a/tools/kangaroo_est.py b/tools/kangaroo_est.py new file mode 100644 index 00000000..33ca025f --- /dev/null +++ b/tools/kangaroo_est.py @@ -0,0 +1,284 @@ +#!/usr/bin/env python3 +"""Runtime / cost / DP-storage estimator for Pollard-kangaroo attacks on the +Bitcoin puzzle addresses that have an exposed public key (every 5th puzzle, +75..160). + +Work model +---------- +Puzzle n has its private key uniformly in [2^(n-1), 2^n), so the interval +width is W = 2^(n-1) and the expected number of secp256k1 group operations is + + ops = K * sqrt(W) = K * 2^((n-1)/2) + +K = 1.15 for the SOTA 3-kangaroo-with-symmetry method (RCKangaroo, +PSCKangaroo, collider). Classic JeanLucPons Kangaroo is K ~= 1.6-2.1 AND is +capped at a 125-bit interval, so it cannot run puzzle 130 and up at all. + +sqrt(W) is the generic-group lower bound (Shoup). There is no known sub-sqrt +attack on secp256k1 ECDLP, so the op count is a floor, not an estimate. The +only lever is operations per dollar. +""" + +import argparse +import math + +K_SOTA = 1.15 # RCKangaroo SOTA (3 kangaroo types + symmetry + loop handling) +K_3WAY = 1.6 # RCKangaroo 3-way +K_JLP = 2.08 # JeanLucPons Kangaroo, per its own README ("2.08*sqrt(k2-k1)") + +SECONDS_PER_YEAR = 365.25 * 24 * 3600 +HOURS_PER_YEAR = 365.25 * 24 + +# (GKeys/s, watts). 4090/5090 are RCKangaroo's own published figures; the rest +# are scaled estimates and are NOT measured. +# CALIBRATE. RCKangaroo prints real GKeys/s on startup -- clocks, driver, +# cooling and the chosen -dp move it 20%+. Paper numbers lie. +GPUS = { + "rtx3090": (6.0, 350), # estimated + "rtx4080": (9.0, 320), # estimated + "rtx4090": (14.5, 450), # measured, RCKangaroo README + "rtx5090": (19.3, 575), # measured, RCKangaroo README (turbo kernels) + "a100": (8.0, 400), # estimated + "h100": (15.0, 700), # estimated +} + +# JeanLucPons measured the DP overhead against r = nbKangaroo * 2^dp / sqrt(N): +# r=4.0 -> 71% wasted, r=0.125 -> 4% wasted. overhead ~= 0.18*r fits that. +# Using the measured curve, not a hand-waved fraction of total work. +DP_OVERHEAD_SLOPE = 0.18 + + +def expected_ops(n, k=K_SOTA): + """Expected group operations to solve puzzle n.""" + return k * 2.0 ** ((n - 1) / 2.0) + + +# Already swept: 1..70 plus every 5th from 75 to 130. Their coins are gone, so +# the prize column is the schedule value, not money you can win. +SOLVED = set(range(1, 71)) | set(range(75, 131, 5)) + + +def prize_btc(n): + """Reward held by puzzle n after the 2023 top-up (0.1 * n BTC for 51..160).""" + return round(0.1 * n, 1) + + +def dp_overhead(n, kangaroos, d): + """Fraction of work wasted on DP tails, per JeanLucPons' measured curve.""" + r = kangaroos * 2.0 ** d / 2.0 ** ((n - 1) / 2.0) + return DP_OVERHEAD_SLOPE * r + + +def dp_window(n, kangaroos, ram_bytes, dp_entry_bytes=16, overhead_frac=0.05, + k=K_SOTA): + """Feasible range of DP bits d. + + Lower bound from RAM: stored distinguished points ~= ops / 2^d, so a small + d means a huge table. Upper bound from overhead: each kangaroo walks an + extra ~2^d steps to reach a DP, and JLP's measured curve puts the waste at + ~0.18 * kangaroos * 2^d / sqrt(W). + + Returns (d_min, d_max). d_min > d_max means infeasible as configured -- + more RAM, or a smaller herd. + """ + ops = expected_ops(n, k) + d_min = max(0.0, math.log2(ops * dp_entry_bytes / ram_bytes)) + sqrt_w = 2.0 ** ((n - 1) / 2.0) + d_max = math.log2(overhead_frac * sqrt_w / (DP_OVERHEAD_SLOPE * kangaroos)) + return d_min, d_max + + +def estimate(n, gpu="rtx4090", count=1, gpu_speed=None, watts=None, + usd_per_gpu_hour=0.35, usd_per_kwh=0.10, btc_usd=100_000.0, + ram_gb=64.0, kangaroos_per_gpu=2 ** 21, dp_entry_bytes=16, + k=K_SOTA): + default_speed, default_watts = GPUS.get(gpu, (None, None)) + speed = gpu_speed if gpu_speed is not None else default_speed + w = watts if watts is not None else default_watts + if speed is None or w is None: + missing = " and ".join(x for x, v in (("--speed", speed), ("--watts", w)) + if v is None) + raise SystemExit("gpu %r has no built-in profile -- pass %s" % (gpu, missing)) + + ops = expected_ops(n, k) + rate = speed * 1e9 * count # ops/sec for the whole fleet + seconds = ops / rate + gpu_hours = (ops / (speed * 1e9)) / 3600.0 + + rental = gpu_hours * usd_per_gpu_hour + kwh = gpu_hours * (w / 1000.0) + power = kwh * usd_per_kwh + + prize = prize_btc(n) + value = prize * btc_usd + + d_min, d_max = dp_window(n, kangaroos_per_gpu * count, + ram_gb * 1024 ** 3, dp_entry_bytes, k=k) + + return { + "puzzle": n, + "interval_bits": n - 1, + "ops": ops, + "ops_log2": math.log2(ops), + "fleet_rate": rate, + "seconds": seconds, + "years": seconds / SECONDS_PER_YEAR, + "days": seconds / 86400.0, + "gpu_years": gpu_hours / HOURS_PER_YEAR, + "gpu_hours": gpu_hours, + "rental_usd": rental, + "kwh": kwh, + "power_usd": power, + "prize_btc": prize, + "prize_usd": value, + "roi_rental": value / rental if rental else float("inf"), + "roi_power": value / power if power else float("inf"), + "breakeven_btc_rental": rental / prize, + "breakeven_btc_power": power / prize, + "dp_min": d_min, + "dp_max": d_max, + "dp_feasible": d_min <= d_max, + "dp_entries_at_min": ops / 2 ** d_min, + } + + +def human_time(seconds): + y = seconds / SECONDS_PER_YEAR + if y >= 1: + return "%.4g years" % y + d = seconds / 86400.0 + if d >= 1: + return "%.4g days" % d + return "%.4g hours" % (seconds / 3600.0) + + +def report(r, gpu, count): + solved = r["puzzle"] in SOLVED + print("puzzle #%d interval 2^%d prize %.1f BTC%s" + % (r["puzzle"], r["interval_bits"], r["prize_btc"], + " *** ALREADY SOLVED -- address swept, prize is $0 ***" if solved else "")) + print(" expected ops %.3g (2^%.1f)" % (r["ops"], r["ops_log2"])) + print(" fleet %d x %s = %.4g ops/s" + % (count, gpu, r["fleet_rate"])) + print(" wall clock %s" % human_time(r["seconds"])) + print(" total GPU-years %.4g" % r["gpu_years"]) + print(" rental cost $%s" % f"{r['rental_usd']:,.0f}") + print(" power only $%s (%.3g kWh)" + % (f"{r['power_usd']:,.0f}", r["kwh"])) + print(" prize value $%s at $%s/BTC" + % (f"{r['prize_usd']:,.0f}", f"{r['prize_usd']/r['prize_btc']:,.0f}")) + print(" ROI vs rental %.3gx %s" + % (r["roi_rental"], "PROFIT" if r["roi_rental"] > 1 else "LOSS")) + print(" ROI vs power only %.3gx %s" + % (r["roi_power"], "PROFIT" if r["roi_power"] > 1 else "LOSS")) + print(" break-even BTC $%s rental / $%s power-only" + % (f"{r['breakeven_btc_rental']:,.0f}", + f"{r['breakeven_btc_power']:,.0f}")) + if solved: + print(" NOTE every ROI line above is fiction: this puzzle is") + print(" solved and the coins are already spent.") + if r["dp_feasible"]: + print(" usable -dp %d .. %d (%.3g DP entries at dp=%d)" + % (math.ceil(r["dp_min"]), math.floor(r["dp_max"]), + r["dp_entries_at_min"], math.ceil(r["dp_min"]))) + else: + print(" usable -dp NONE: RAM floor dp>=%.1f exceeds overhead " + "ceiling dp<=%.1f -- add RAM or shrink the herd" + % (r["dp_min"], r["dp_max"])) + # ponytail: point estimate only. Kangaroo runtime has real variance -- a run + # can finish well under or over expected ops. Not modelled; treat as a mean. + + +def compare(puzzles, **kw): + print("%-8s %-10s %-12s %-12s %-14s %-10s" % + ("puzzle", "ops log2", "GPU-years", "rental $", "prize BTC", "ROI")) + for n in puzzles: + r = estimate(n, **kw) + print("%-8d %-10.1f %-12.4g %-12s %-14.1f %-10.3g" + % (n, r["ops_log2"], r["gpu_years"], + f"{r['rental_usd']:,.0f}", r["prize_btc"], r["roi_rental"])) + + +def demo(): + # Published figure for #135: ~1.15 * sqrt(2^134) ~= 1.7e20 ops (2^67.2). + o135 = expected_ops(135) + assert 1.6e20 < o135 < 1.8e20, o135 + assert abs(math.log2(o135) - 67.2) < 0.1 + + # Each +1 on the puzzle number is sqrt(2) more work; +5 is 2^2.5 = 5.657x. + assert abs(expected_ops(140) / expected_ops(135) - 2 ** 2.5) < 1e-9 + assert abs(expected_ops(140) / expected_ops(130) - 2 ** 5) < 1e-9 + + # Prize schedule after the 2023 top-up. + assert prize_btc(66) == 6.6 and prize_btc(140) == 14.0 + # Solved set must cover the swept ranges and exclude the live targets. + assert 66 in SOLVED and 120 in SOLVED and 130 in SOLVED + assert 135 not in SOLVED and 140 not in SOLVED and 71 not in SOLVED + + # #130 on a single 4090 (14.5 GH/s) is ~65 GPU-years; #140 is 32x that. + r130 = estimate(130) + r140 = estimate(140) + assert 55 < r130["gpu_years"] < 80, r130["gpu_years"] + assert abs(r140["gpu_years"] / r130["gpu_years"] - 32) < 1e-6 + + # JLP's K=2.08 costs 1.81x more work than SOTA's K=1.15 -- matches + # RCKangaroo's "1.8 times less required operations" claim. + assert abs(expected_ops(140, K_JLP) / expected_ops(140, K_SOTA) - 1.809) < 0.01 + + # JLP's own measured DP overhead: r=4.0 -> ~71% wasted. + sqrt_w = 2.0 ** ((140 - 1) / 2.0) + kang = 4.0 * sqrt_w / 2.0 ** 30 + assert abs(dp_overhead(140, kang, 30) - 0.72) < 0.02 + + # Fleet size divides wall clock but never total GPU-hours. + solo = estimate(140, count=1) + fleet = estimate(140, count=1000) + assert abs(solo["gpu_hours"] - fleet["gpu_hours"]) < 1e-3 + assert abs(solo["seconds"] / fleet["seconds"] - 1000) < 1e-6 + + # DP window: more RAM lowers the floor, a bigger herd lowers the ceiling. + lo, hi = dp_window(140, 2 ** 21, 64 * 1024 ** 3) + lo_more_ram, _ = dp_window(140, 2 ** 21, 512 * 1024 ** 3) + _, hi_big_herd = dp_window(140, 2 ** 30, 64 * 1024 ** 3) + assert lo_more_ram < lo and hi_big_herd < hi + assert lo <= hi, "expected 64GB/2^21 to be feasible at #140" + + print("self-check OK") + + +if __name__ == "__main__": + p = argparse.ArgumentParser(description=__doc__, + formatter_class=argparse.RawDescriptionHelpFormatter) + p.add_argument("puzzle", nargs="?", type=int, default=140) + p.add_argument("--gpu", default="rtx4090", choices=sorted(GPUS) + ["custom"]) + p.add_argument("--count", type=int, default=1, help="number of GPUs") + p.add_argument("--speed", type=float, help="GKeys/s per GPU (overrides preset)") + p.add_argument("--watts", type=float, help="watts per GPU (overrides preset)") + p.add_argument("--rate", type=float, default=0.35, help="USD per GPU-hour") + p.add_argument("--kwh", type=float, default=0.10, help="USD per kWh") + p.add_argument("--btc", type=float, default=100_000.0, help="assumed BTC price") + p.add_argument("--ram", type=float, default=64.0, help="host RAM in GB for DPs") + p.add_argument("--kangaroos", type=int, default=2 ** 21, + help="concurrent kangaroos per GPU (see solver output)") + p.add_argument("--dp-bytes", type=int, default=16, + help="bytes per stored DP (16 = PSCKangaroo compact, 32 = RCKangaroo)") + p.add_argument("--k", type=float, default=K_SOTA, + help="method constant: 1.15 SOTA, 1.6 3-way, 2.08 JeanLucPons") + p.add_argument("--compare", action="store_true", + help="table across 130/135/140/145/150") + p.add_argument("--self-check", action="store_true") + a = p.parse_args() + + if a.self_check: + demo() + raise SystemExit(0) + + kw = dict(gpu=a.gpu, count=a.count, gpu_speed=a.speed, watts=a.watts, + usd_per_gpu_hour=a.rate, usd_per_kwh=a.kwh, btc_usd=a.btc, + ram_gb=a.ram, kangaroos_per_gpu=a.kangaroos, dp_entry_bytes=a.dp_bytes, + k=a.k) + + if a.compare: + compare([130, 135, 140, 145, 150], **kw) + else: + report(estimate(a.puzzle, **kw), a.gpu, a.count) diff --git a/tools/verify_target.py b/tools/verify_target.py new file mode 100644 index 00000000..388216d7 --- /dev/null +++ b/tools/verify_target.py @@ -0,0 +1,152 @@ +#!/usr/bin/env python3 +"""Verify a puzzle target before spending GPU time on it. + +Derives the P2PKH address from the compressed public key and checks it against +the published puzzle address, and checks the range is the canonical +[2^(n-1), 2^n - 1]. A wrong pubkey pasted from a forum makes an entire run +worthless, and nothing in the solver would tell you. + +Pure stdlib: secp256k1 point decompression + hash160 + base58check. +""" + +import hashlib +import sys + +P = 2**256 - 2**32 - 977 +N = 0xFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFEBAAEDCE6AF48A03BBFD25E8CD0364141 +B58 = "123456789ABCDEFGHJKLMNPQRSTUVWXYZabcdefghijkmnopqrstuvwxyz" + + +def decompress(pub_hex): + """Compressed SEC pubkey -> (x, y), verified on-curve.""" + raw = bytes.fromhex(pub_hex) + if len(raw) != 33 or raw[0] not in (2, 3): + raise ValueError("not a 33-byte compressed pubkey") + x = int.from_bytes(raw[1:], "big") + if x >= P: + raise ValueError("x out of field") + y = pow((x * x * x + 7) % P, (P + 1) // 4, P) # P % 4 == 3 + if (y * y - (x * x * x + 7)) % P != 0: + raise ValueError("x is not on the curve") + if y % 2 != raw[0] % 2: + y = P - y + return x, y + + +def hash160(b): + return hashlib.new("ripemd160", hashlib.sha256(b).digest()).digest() + + +def b58check(payload): + chk = hashlib.sha256(hashlib.sha256(payload).digest()).digest()[:4] + n = int.from_bytes(payload + chk, "big") + out = "" + while n: + n, r = divmod(n, 58) + out = B58[r] + out + return "1" * (len(payload + chk) - len((payload + chk).lstrip(b"\0"))) + out + + +def address(pub_hex): + return b58check(b"\x00" + hash160(bytes.fromhex(pub_hex))) + + +def check(puzzle, pub_hex, addr, start_hex, stop_hex): + ok = True + x, y = decompress(pub_hex) # raises if off-curve + + derived = address(pub_hex) + match = derived == addr + ok &= match + print("puzzle #%d" % puzzle) + print(" pubkey on curve yes") + print(" address expected %s" % addr) + print(" address derived %s %s" % (derived, "MATCH" if match else "MISMATCH")) + + start, stop = int(start_hex, 16), int(stop_hex, 16) + want_start, want_stop = 2 ** (puzzle - 1), 2 ** puzzle - 1 + rng = (start == want_start and stop == want_stop) + ok &= rng + print(" range 2^%d .. 2^%d-1 %s" + % (puzzle - 1, puzzle, "OK" if rng else "UNEXPECTED")) + print(" interval bits %d" % (puzzle - 1)) + if stop >= N: + print(" WARNING: range exceeds the curve order") + ok = False + return ok + + +def demo(): + # Puzzle #105, solved, key published in the repo's puzzle32.txt. If the + # address derivation is right, this key's pubkey must yield #105's address. + priv = 0x16F14FC2054CD87EE6396B33DF3 + assert 2**104 <= priv < 2**105 + + # k*G by double-and-add, to get the pubkey from the known private key. + GX = 0x79BE667EF9DCBBAC55A06295CE870B07029BFCDB2DCE28D959F2815B16F81798 + GY = 0x483ADA7726A3C4655DA4FBFC0E1108A8FD17B448A68554199C47D08FFB10D4B8 + + def add(p, q): + if p is None: + return q + if q is None: + return p + if p[0] == q[0]: + if (p[1] + q[1]) % P == 0: + return None + l = 3 * p[0] * p[0] * pow(2 * p[1], -1, P) % P + else: + l = (q[1] - p[1]) * pow(q[0] - p[0], -1, P) % P + rx = (l * l - p[0] - q[0]) % P + return rx, (l * (p[0] - rx) - p[1]) % P + + r, base, k = None, (GX, GY), priv + while k: + if k & 1: + r = add(r, base) + base = add(base, base) + k >>= 1 + pub = "%02x%064x" % (2 + (r[1] & 1), r[0]) + assert address(pub) == "1CMjscKB3QW7SDyQ4c3C3DEUHiHRhiZVib", address(pub) + + # Round-trip: decompressing our own encoding returns the same point. + assert decompress(pub) == r + + # An off-curve x must be rejected, not silently accepted. Find one honestly: + # x is on the curve iff x^3+7 is a quadratic residue mod p. + bad_x = next(x for x in range(1, 200) + if pow((x * x * x + 7) % P, (P - 1) // 2, P) != 1) + try: + decompress("02" + "%064x" % bad_x) + raise AssertionError("x=%d is off-curve; should have raised" % bad_x) + except ValueError: + pass + + print("verify_target self-check OK (derives #105 from its published key)") + + +TARGETS = { + 135: ("02145D2611C823A396EF6712CE0F712F09B9B4F3135E3E0AA3230FB9B6D08D1E16", + "16RGFo6hjq9ym6Pj7N5H7L1NR1rVPJyw2v", + "4000000000000000000000000000000000", + "7FFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFF"), + 120: ("02CEB6CBBCDBDF5EF7150682150F4CE2C6F4807B349827DCDBDD1F2EFA885A2630", + "17s2b9ksz5y7abUm92cHwG8jEPCzK3dLnT", + "800000000000000000000000000000", + "FFFFFFFFFFFFFFFFFFFFFFFFFFFFFF"), + 140: ("031F6A332D3C5C4F2DE2378C012F429CD109BA07D69690C6C701B6BB87860D6640", + "1QKBaU6WAeycb3DbKbLBkX7vJiaS8r42Xo", + "80000000000000000000000000000000000", + "FFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFF"), +} + +if __name__ == "__main__": + if len(sys.argv) > 1 and sys.argv[1] == "--self-check": + demo() + raise SystemExit(0) + bad = 0 + for n in (int(a) for a in (sys.argv[1:] or ["135", "140"])): + if not check(n, *TARGETS[n]): + bad += 1 + print() + raise SystemExit(1 if bad else 0) From 4276449905ea817403c7abeac4e9e370c8d8f3a0 Mon Sep 17 00:00:00 2001 From: Dileep-Kumar-5 Date: Thu, 6 Aug 2026 06:27:48 +0530 Subject: [PATCH 6/7] fix: teach _ConvertSMVer2Cores about Ampere and newer The table stops at sm_75, so every GPU from Ampere on falls through to the `return 0` default. Two consequences, and only the first is cosmetic: - the startup banner prints "(16x0 cores)" on an RTX 3050 - GetGridSize() computes `*y = 2 * _ConvertSMVer2Cores(...)`, gets 0, and lands on the `if(*y <= 0) *y = 128` fallback -- so the default grid Y is 128 instead of the intended 2 x 128 = 256 on any sm_86 card Added sm_80/86/87/89/90/a0/a1/c0 with NVIDIA's per-SM core counts. Measured on an RTX 3050 Laptop, same binary, -g forced, back to back: grid 32x128 (old default) 338.9 MK/s grid 32x256 (new default) 345.5 MK/s +2.0% Modest. A short sample suggested +16%, but that did not survive a matched window -- this laptop GPU throttles, so only back-to-back deltas mean anything. The grid also doubles kangaroo memory, 57 -> 105 MB. Nothing here depends on the 256-bit distance work; it is a correctness fix to a lookup table that has simply gone stale. Co-Authored-By: Claude Opus 5 --- GPU/GPUEngine.cu | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/GPU/GPUEngine.cu b/GPU/GPUEngine.cu index cb6fcba1..8b61e472 100644 --- a/GPU/GPUEngine.cu +++ b/GPU/GPUEngine.cu @@ -121,6 +121,14 @@ int _ConvertSMVer2Cores(int major,int minor) { { 0x70, 64 }, { 0x72, 64 }, { 0x75, 64 }, + { 0x80, 64 }, // Ampere GA100 + { 0x86, 128 }, // Ampere GA10x + { 0x87, 128 }, // Ampere Orin + { 0x89, 128 }, // Ada Lovelace + { 0x90, 128 }, // Hopper + { 0xa0, 128 }, // Blackwell GB100 + { 0xa1, 128 }, + { 0xc0, 128 }, // Blackwell GB20x { -1, -1 } }; int index = 0; From af9c5aa4b745b23785cb0d1e8ce2f34249fbf9f4 Mon Sep 17 00:00:00 2001 From: Dileep-Kumar-5 Date: Thu, 6 Aug 2026 06:59:26 +0530 Subject: [PATCH 7/7] feat: make the distance width a compile-time option (default = upstream) The 256-bit distance is now opt-in. DIST_WORDS selects the width: 2 by default, 4 with -DWIDE_DIST (or by uncommenting WIDE_DIST in Constants.h). This addresses the main objection to the previous form. A default build now reproduces upstream's layout exactly -- ENTRY 32 bytes, ITEM 56, KSIZE 10, DP packet 40, kangaroo transfer 16, and the original work-file and protocol magics -- so it reads and writes upstream's files and talks to upstream peers. Only -DWIDE_DIST changes any of that. Rather than #ifdef every use site, a dist_t typedef plus DIST_WORDS lets one implementation serve both widths: the sign and kangaroo-type flags always occupy the top two bits of the top word, so Convert, CalcDistAndType, CreateEntry, sameDist and the four network checksum accumulations became loops over DIST_WORDS. That also removed the legacy 126-bit Add/Convert/CalcDistAndType overloads and Widen() -- at DIST_WORDS == 2 dist_t simply is int128_t. Only genuinely width-specific code is conditional: the two extra device slots in Load/StoreKangaroos, the four extra uint32 in OutputDP, the Add128/Add256 selection behind AddDist, and the format magics. Exceeding the compiled limit still aborts, and now names the fix: HashTable::Convert: travelled distance exceeds 126 bits. Interval is too large for the DP entry format (max 125 bits). Rebuild with -DWIDE_DIST for intervals up to 253 bits. Aborting rather than storing a truncated distance. Verified, both builds, RTX 3050 Laptop / CUDA 13.3 / sm_86: - default and wide both solve the 56-bit sample key correctly - default on #120 (2^119): 34,906 DPs, 100.000% -wcheck - default on #140 (2^139): aborts as above, exit 255 - wide on #140 (2^139): 34,762 DPs, 100.000% -wcheck - default writes magic 0xFA6A8001, i.e. upstream's HEADW - each build rejects the other's work file ("Not a work file") tools/check_layout.py and tools/test_distfield.py now verify BOTH widths in one run, and assert the default reproduces upstream's sizes. Neither needs a compiler or a GPU. build_gpu.bat takes an optional "wide" arg. README documents the option, the costs and the incompatibility. Co-Authored-By: Claude Opus 5 --- Backup.cpp | 10 +- Constants.h | 21 +++ GPU/GPUCompute.h | 4 +- GPU/GPUEngine.cu | 19 ++- GPU/GPUEngine.h | 15 +- GPU/GPUMath.h | 52 ++++-- HashTable.cpp | 61 +++---- HashTable.h | 45 ++--- Kangaroo.cpp | 2 +- Kangaroo.h | 24 ++- Network.cpp | 77 ++++----- README.md | 55 +++++- build_gpu.bat | 24 ++- tools/check_layout.py | 361 ++++++++++++++++++++-------------------- tools/test_distfield.py | 159 +++++++++--------- 15 files changed, 531 insertions(+), 398 deletions(-) diff --git a/Backup.cpp b/Backup.cpp index 71db105b..9fb0b0f4 100644 --- a/Backup.cpp +++ b/Backup.cpp @@ -230,7 +230,7 @@ void Kangaroo::FetchWalks(uint64_t nbWalk,Int *x,Int *y,Int *d) { } -void Kangaroo::FetchWalks(uint64_t nbWalk,std::vector& kangs,Int* x,Int* y,Int* d) { +void Kangaroo::FetchWalks(uint64_t nbWalk,std::vector& kangs,Int* x,Int* y,Int* d) { uint64_t n = 0; @@ -293,7 +293,7 @@ void Kangaroo::FectchKangaroos(TH_PARAM *threads) { double sFetch = Timer::get_tick(); // From server - vector kangs; + vector kangs; if(saveKangarooByServer) { ::printf("FectchKangaroosFromServer"); if(!GetKangaroosFromServer(workFile,kangs)) @@ -492,14 +492,14 @@ void Kangaroo::SaveWork(uint64_t totalCount,double totalTime,TH_PARAM *threads,i if(saveKangarooByServer) { ::printf("\nSaveWork (Kangaroo->Server): %s",fileName.c_str()); - vector kangs; + vector kangs; for(int i = 0; i < nbThread; i++) totalWalk += threads[i].nbKangaroo; kangs.reserve(totalWalk); for(int i = 0; i < nbThread; i++) { int128_t X; - int256_t D; + dist_t D; uint64_t h; for(uint64_t n = 0; n < threads[i].nbKangaroo; n++) { HashTable::Convert(&threads[i].px[n],&threads[i].distance[n],n%2,&h,&X,&D); @@ -507,7 +507,7 @@ void Kangaroo::SaveWork(uint64_t totalCount,double totalTime,TH_PARAM *threads,i } } SendKangaroosToServer(fileName,kangs); - size = kangs.size()*32 + 16; + size = kangs.size()*(8*DIST_WORDS) + 16; goto end; } else { diff --git a/Constants.h b/Constants.h index 25ae0434..36de4377 100644 --- a/Constants.h +++ b/Constants.h @@ -24,6 +24,27 @@ // Use symmetry //#define USE_SYMMETRY +// Travelled-distance field width. +// +// Default (off): the original layout -- 128 bits per distance, of which +// b127=sign, b126=kangaroo type and b125..b0 the magnitude. That caps the +// search interval at 125 bits, and exceeding it used to corrupt keys +// silently; it now aborts. +// +// Uncomment, or build with -DWIDE_DIST, for a 256-bit field (254-bit +// magnitude, intervals up to 253 bits). Needed for puzzle 130 and above. +// Costs: DP tables +50% RAM, device kangaroo memory +20%, DP packets +40%, +// kangaroo transfers +100%, and roughly 4-7% throughput. Work files and the +// client/server protocol are NOT interchangeable between the two builds -- +// the format magics differ so a mismatch is refused, not misparsed. +//#define WIDE_DIST + +#ifdef WIDE_DIST +#define DIST_WORDS 4 +#else +#define DIST_WORDS 2 +#endif + // Number of random jumps // Max 512 for the GPU #define NB_JUMP 32 diff --git a/GPU/GPUCompute.h b/GPU/GPUCompute.h index ff03fe0a..ea15e6f0 100644 --- a/GPU/GPUCompute.h +++ b/GPU/GPUCompute.h @@ -23,7 +23,7 @@ __device__ void ComputeKangaroos(uint64_t *kangaroos,uint32_t maxFound,uint32_t uint64_t px[GPU_GRP_SIZE][4]; uint64_t py[GPU_GRP_SIZE][4]; - uint64_t dist[GPU_GRP_SIZE][4]; + uint64_t dist[GPU_GRP_SIZE][DIST_WORDS]; #ifdef USE_SYMMETRY uint64_t lastJump[GPU_GRP_SIZE]; #endif @@ -86,7 +86,7 @@ __device__ void ComputeKangaroos(uint64_t *kangaroos,uint32_t maxFound,uint32_t Load256(px[g],rx); Load256(py[g],ry); - Add256(dist[g],jD[jmp]); + AddDist(dist[g],jD[jmp]); #ifdef USE_SYMMETRY if(ModPositive256(py[g])) diff --git a/GPU/GPUEngine.cu b/GPU/GPUEngine.cu index 8b61e472..e27b64c4 100644 --- a/GPU/GPUEngine.cu +++ b/GPU/GPUEngine.cu @@ -424,12 +424,14 @@ void GPUEngine::SetKangaroos(Int *px,Int *py,Int *d) { if(idx % 2 == WILD) dOff.ModAddK1order(&wildOffset); inputKangarooPinned[g * strideSize + t + 8 * nbThreadPerGroup] = dOff.bits64[0]; inputKangarooPinned[g * strideSize + t + 9 * nbThreadPerGroup] = dOff.bits64[1]; +#if DIST_WORDS == 4 inputKangarooPinned[g * strideSize + t + 10 * nbThreadPerGroup] = dOff.bits64[2]; inputKangarooPinned[g * strideSize + t + 11 * nbThreadPerGroup] = dOff.bits64[3]; +#endif #ifdef USE_SYMMETRY // Last jump - inputKangarooPinned[t + 12 * nbThreadPerGroup] = (uint64_t)NB_JUMP; + inputKangarooPinned[t + (8 + DIST_WORDS) * nbThreadPerGroup] = (uint64_t)NB_JUMP; #endif idx++; @@ -491,8 +493,10 @@ void GPUEngine::GetKangaroos(Int *px,Int *py,Int *d) { dOff.SetInt32(0); dOff.bits64[0] = inputKangarooPinned[g * strideSize + t + 8 * nbThreadPerGroup]; dOff.bits64[1] = inputKangarooPinned[g * strideSize + t + 9 * nbThreadPerGroup]; +#if DIST_WORDS == 4 dOff.bits64[2] = inputKangarooPinned[g * strideSize + t + 10 * nbThreadPerGroup]; dOff.bits64[3] = inputKangarooPinned[g * strideSize + t + 11 * nbThreadPerGroup]; +#endif if(idx % 2 == WILD) dOff.ModSubK1order(&wildOffset); d[idx].Set(&dOff); @@ -547,10 +551,12 @@ void GPUEngine::SetKangaroo(uint64_t kIdx,Int *px,Int *py,Int *d) { cudaMemcpy(inputKangaroo + (b * blockSize + g * strideSize + t + 8 * nbThreadPerGroup),inputKangarooPinned,8,cudaMemcpyHostToDevice); inputKangarooPinned[0] = dOff.bits64[1]; cudaMemcpy(inputKangaroo + (b * blockSize + g * strideSize + t + 9 * nbThreadPerGroup),inputKangarooPinned,8,cudaMemcpyHostToDevice); +#if DIST_WORDS == 4 inputKangarooPinned[0] = dOff.bits64[2]; cudaMemcpy(inputKangaroo + (b * blockSize + g * strideSize + t + 10 * nbThreadPerGroup),inputKangarooPinned,8,cudaMemcpyHostToDevice); inputKangarooPinned[0] = dOff.bits64[3]; cudaMemcpy(inputKangaroo + (b * blockSize + g * strideSize + t + 11 * nbThreadPerGroup),inputKangarooPinned,8,cudaMemcpyHostToDevice); +#endif #ifdef USE_SYMMETRY // Last jump @@ -584,8 +590,8 @@ void GPUEngine::SetParams(uint64_t dpMask,Int *distance,Int *px,Int *py) { this->dpMask = dpMask; for(int i=0;i< NB_JUMP;i++) - memcpy(jumpPinned + 4*i,distance[i].bits64,32); - cudaMemcpyToSymbol(jD,jumpPinned,jumpSize); + memcpy(jumpPinned + DIST_WORDS*i,distance[i].bits64,8*DIST_WORDS); + cudaMemcpyToSymbol(jD,jumpPinned,NB_JUMP*8*DIST_WORDS); cudaError_t err = cudaGetLastError(); if(err != cudaSuccess) { printf("GPUEngine: SetParams: Failed to copy to constant memory: %s\n",cudaGetErrorString(err)); @@ -677,7 +683,7 @@ bool GPUEngine::Launch(std::vector &hashFound,bool spinWait) { uint32_t *itemPtr = outputItemPinned + (i*ITEM_SIZE32 + 1); ITEM it; - it.kIdx = *((uint64_t*)(itemPtr + 16)); + it.kIdx = *((uint64_t*)(itemPtr + 8 + 2*DIST_WORDS)); uint64_t *x = (uint64_t *)itemPtr; it.x.bits64[0] = x[0]; @@ -689,8 +695,13 @@ bool GPUEngine::Launch(std::vector &hashFound,bool spinWait) { uint64_t *d = (uint64_t *)(itemPtr + 8); it.d.bits64[0] = d[0]; it.d.bits64[1] = d[1]; +#if DIST_WORDS == 4 it.d.bits64[2] = d[2]; it.d.bits64[3] = d[3]; +#else + it.d.bits64[2] = 0; + it.d.bits64[3] = 0; +#endif it.d.bits64[4] = 0; if(it.kIdx % 2 == WILD) it.d.ModSubK1order(&wildOffset); diff --git a/GPU/GPUEngine.h b/GPU/GPUEngine.h index fe1d00a7..6dcb3b7f 100644 --- a/GPU/GPUEngine.h +++ b/GPU/GPUEngine.h @@ -22,17 +22,18 @@ #include "../Constants.h" #include "../SECPK1/SECP256k1.h" -// Words per kangaroo in device memory: px[4] + py[4] + dist[4] (+ lastJump). -// dist was 2 words (126bit distance, 125bit interval cap); it is 4 words now. +// Words per kangaroo in device memory: px[4] + py[4] + dist[DIST_WORDS] +// (+ lastJump). DIST_WORDS is 2 by default (126bit distance, 125bit interval +// cap) and 4 with WIDE_DIST. #ifdef USE_SYMMETRY -#define KSIZE 13 +#define KSIZE (9 + DIST_WORDS) #else -#define KSIZE 12 +#define KSIZE (8 + DIST_WORDS) #endif -// x[8] + d[8] + kIdx[2], in uint32 -#define ITEM_SIZE 72 -#define ITEM_SIZE32 (ITEM_SIZE/4) +// x[8] + d[2*DIST_WORDS] + kIdx[2], in uint32 +#define ITEM_SIZE32 (8 + 2*DIST_WORDS + 2) +#define ITEM_SIZE (ITEM_SIZE32*4) typedef struct { Int x; diff --git a/GPU/GPUMath.h b/GPU/GPUMath.h index ae008c31..26ff4170 100644 --- a/GPU/GPUMath.h +++ b/GPU/GPUMath.h @@ -48,7 +48,7 @@ #define MADDS(r,a,b,c) asm volatile ("madc.hi.s64 %0, %1, %2, %3;" : "=l"(r) : "l"(a), "l"(b), "l"(c)); // Jump distance -__device__ __constant__ uint64_t jD[NB_JUMP][4]; +__device__ __constant__ uint64_t jD[NB_JUMP][DIST_WORDS]; // jump points __device__ __constant__ uint64_t jPx[NB_JUMP][4]; __device__ __constant__ uint64_t jPy[NB_JUMP][4]; @@ -131,6 +131,13 @@ __device__ __constant__ uint64_t _O[] = { 0xBFD25E8CD0364141ULL,0xBAAEDCE6AF48A0 UADDC1((r)[2], (a)[2]); \ UADD1((r)[3], (a)[3]);} +// Distance accumulate at whichever width is compiled in. +#if DIST_WORDS == 4 +#define AddDist(r,a) Add256(r,a) +#else +#define AddDist(r,a) Add128(r,a) +#endif + // --------------------------------------------------------------------------------------- #define Neg(r) {\ @@ -181,6 +188,16 @@ USUB(r[4],0ULL,r[4]); } // --------------------------------------------------------------------------------------- +#if DIST_WORDS == 4 +#define DP_EXTRA_DIST_WORDS(d) \ +out[pos*ITEM_SIZE32 + 13] = ((uint32_t *)d)[4]; \ +out[pos*ITEM_SIZE32 + 14] = ((uint32_t *)d)[5]; \ +out[pos*ITEM_SIZE32 + 15] = ((uint32_t *)d)[6]; \ +out[pos*ITEM_SIZE32 + 16] = ((uint32_t *)d)[7]; +#else +#define DP_EXTRA_DIST_WORDS(d) +#endif + #define OutputDP(x,d,idx) {\ out[pos*ITEM_SIZE32 + 1] = ((uint32_t *)x)[0]; \ out[pos*ITEM_SIZE32 + 2] = ((uint32_t *)x)[1]; \ @@ -194,20 +211,17 @@ out[pos*ITEM_SIZE32 + 9] = ((uint32_t *)d)[0]; \ out[pos*ITEM_SIZE32 + 10] = ((uint32_t *)d)[1]; \ out[pos*ITEM_SIZE32 + 11] = ((uint32_t *)d)[2]; \ out[pos*ITEM_SIZE32 + 12] = ((uint32_t *)d)[3]; \ -out[pos*ITEM_SIZE32 + 13] = ((uint32_t *)d)[4]; \ -out[pos*ITEM_SIZE32 + 14] = ((uint32_t *)d)[5]; \ -out[pos*ITEM_SIZE32 + 15] = ((uint32_t *)d)[6]; \ -out[pos*ITEM_SIZE32 + 16] = ((uint32_t *)d)[7]; \ -out[pos*ITEM_SIZE32 + 17] = ((uint32_t *)idx)[0]; \ -out[pos*ITEM_SIZE32 + 18] = ((uint32_t *)idx)[1]; \ +DP_EXTRA_DIST_WORDS(d) \ +out[pos*ITEM_SIZE32 + 2*DIST_WORDS + 9] = ((uint32_t *)idx)[0]; \ +out[pos*ITEM_SIZE32 + 2*DIST_WORDS + 10] = ((uint32_t *)idx)[1]; \ } // --------------------------------------------------------------------------------------- #ifdef USE_SYMMETRY -__device__ void LoadKangaroos(uint64_t *a,uint64_t px[GPU_GRP_SIZE][4],uint64_t py[GPU_GRP_SIZE][4],uint64_t dist[GPU_GRP_SIZE][4],uint64_t *jumps) { +__device__ void LoadKangaroos(uint64_t *a,uint64_t px[GPU_GRP_SIZE][4],uint64_t py[GPU_GRP_SIZE][4],uint64_t dist[GPU_GRP_SIZE][DIST_WORDS],uint64_t *jumps) { #else -__device__ void LoadKangaroos(uint64_t * a,uint64_t px[GPU_GRP_SIZE][4],uint64_t py[GPU_GRP_SIZE][4],uint64_t dist[GPU_GRP_SIZE][4]) { +__device__ void LoadKangaroos(uint64_t * a,uint64_t px[GPU_GRP_SIZE][4],uint64_t py[GPU_GRP_SIZE][4],uint64_t dist[GPU_GRP_SIZE][DIST_WORDS]) { #endif __syncthreads(); @@ -231,17 +245,19 @@ __device__ void LoadKangaroos(uint64_t * a,uint64_t px[GPU_GRP_SIZE][4],uint64_t d64[0] = (a)[IDX + 8 * blockDim.x + stride]; d64[1] = (a)[IDX + 9 * blockDim.x + stride]; +#if DIST_WORDS == 4 d64[2] = (a)[IDX + 10 * blockDim.x + stride]; d64[3] = (a)[IDX + 11 * blockDim.x + stride]; +#endif #ifdef USE_SYMMETRY - jumps[g] = (a)[IDX + 12 * blockDim.x + stride]; + jumps[g] = (a)[IDX + (8 + DIST_WORDS) * blockDim.x + stride]; #endif } } -__device__ void LoadDists(uint64_t* a,uint64_t dist[GPU_GRP_SIZE][4]) { +__device__ void LoadDists(uint64_t* a,uint64_t dist[GPU_GRP_SIZE][DIST_WORDS]) { __syncthreads(); @@ -252,8 +268,10 @@ __device__ void LoadDists(uint64_t* a,uint64_t dist[GPU_GRP_SIZE][4]) { d64[0] = (a)[IDX + 8 * blockDim.x + stride]; d64[1] = (a)[IDX + 9 * blockDim.x + stride]; +#if DIST_WORDS == 4 d64[2] = (a)[IDX + 10 * blockDim.x + stride]; d64[3] = (a)[IDX + 11 * blockDim.x + stride]; +#endif } @@ -290,9 +308,9 @@ __device__ void LoadKangaroo(uint64_t* a,uint32_t stride,uint64_t px[4]) { // --------------------------------------------------------------------------------------- #ifdef USE_SYMMETRY -__device__ void StoreKangaroos(uint64_t *a,uint64_t px[GPU_GRP_SIZE][4],uint64_t py[GPU_GRP_SIZE][4],uint64_t dist[GPU_GRP_SIZE][4],uint64_t *jumps) { +__device__ void StoreKangaroos(uint64_t *a,uint64_t px[GPU_GRP_SIZE][4],uint64_t py[GPU_GRP_SIZE][4],uint64_t dist[GPU_GRP_SIZE][DIST_WORDS],uint64_t *jumps) { #else -__device__ void StoreKangaroos(uint64_t * a,uint64_t px[GPU_GRP_SIZE][4],uint64_t py[GPU_GRP_SIZE][4],uint64_t dist[GPU_GRP_SIZE][4]) { +__device__ void StoreKangaroos(uint64_t * a,uint64_t px[GPU_GRP_SIZE][4],uint64_t py[GPU_GRP_SIZE][4],uint64_t dist[GPU_GRP_SIZE][DIST_WORDS]) { #endif __syncthreads(); @@ -315,11 +333,13 @@ __device__ void StoreKangaroos(uint64_t * a,uint64_t px[GPU_GRP_SIZE][4],uint64_ (a)[IDX + 8 * blockDim.x + stride] = d64[0]; (a)[IDX + 9 * blockDim.x + stride] = d64[1]; +#if DIST_WORDS == 4 (a)[IDX + 10 * blockDim.x + stride] = d64[2]; (a)[IDX + 11 * blockDim.x + stride] = d64[3]; +#endif #ifdef USE_SYMMETRY - (a)[IDX + 12 * blockDim.x + stride] = jumps[g]; + (a)[IDX + (8 + DIST_WORDS) * blockDim.x + stride] = jumps[g]; #endif } @@ -342,7 +362,7 @@ __device__ void StoreKangaroo(uint64_t* a,uint32_t stride,uint64_t px[4],uint64_ } -__device__ void StoreDists(uint64_t* a,uint64_t dist[GPU_GRP_SIZE][4]) { +__device__ void StoreDists(uint64_t* a,uint64_t dist[GPU_GRP_SIZE][DIST_WORDS]) { __syncthreads(); @@ -352,8 +372,10 @@ __device__ void StoreDists(uint64_t* a,uint64_t dist[GPU_GRP_SIZE][4]) { (a)[IDX + 8 * blockDim.x + stride] = d64[0]; (a)[IDX + 9 * blockDim.x + stride] = d64[1]; +#if DIST_WORDS == 4 (a)[IDX + 10 * blockDim.x + stride] = d64[2]; (a)[IDX + 11 * blockDim.x + stride] = d64[3]; +#endif } diff --git a/HashTable.cpp b/HashTable.cpp index 51c97969..03d976fe 100644 --- a/HashTable.cpp +++ b/HashTable.cpp @@ -55,23 +55,22 @@ uint64_t HashTable::GetNbItem() { } -ENTRY *HashTable::CreateEntry(int128_t *x,int256_t *d) { +ENTRY *HashTable::CreateEntry(int128_t *x,dist_t *d) { ENTRY *e = (ENTRY *)malloc(sizeof(ENTRY)); e->x.i64[0] = x->i64[0]; e->x.i64[1] = x->i64[1]; - e->d.i64[0] = d->i64[0]; - e->d.i64[1] = d->i64[1]; - e->d.i64[2] = d->i64[2]; - e->d.i64[3] = d->i64[3]; + for(int i = 0; i < DIST_WORDS; i++) + e->d.i64[i] = d->i64[i]; return e; } -bool HashTable::sameDist(int256_t *a,int256_t *b) { +bool HashTable::sameDist(dist_t *a,dist_t *b) { - return (a->i64[0] == b->i64[0]) && (a->i64[1] == b->i64[1]) && - (a->i64[2] == b->i64[2]) && (a->i64[3] == b->i64[3]); + for(int i = 0; i < DIST_WORDS; i++) + if(a->i64[i] != b->i64[i]) return false; + return true; } @@ -82,7 +81,7 @@ bool HashTable::sameDist(int256_t *a,int256_t *b) { E[h].items[st] = entry; \ E[h].nbItem++;} -void HashTable::Convert(Int *x,Int *d,uint32_t type,uint64_t *h,int128_t *X,int256_t *D) { +void HashTable::Convert(Int *x,Int *d,uint32_t type,uint64_t *h,int128_t *X,dist_t *D) { uint64_t sign = 0; uint64_t type64 = (uint64_t)type << 62; @@ -99,21 +98,27 @@ void HashTable::Convert(Int *x,Int *d,uint32_t type,uint64_t *h,int128_t *X,int2 sign = DIST_SIGN_MASK; } - // The magnitude has to fit in b253..b0. The original code masked the excess - // away, which produced a valid-looking entry and, on collision, a WRONG - // private key with no error. Refuse loudly instead. - if(N.bits64[3] & 0xC000000000000000ULL) { + // The magnitude has to fit in DIST_MAG_BITS. The original code masked the + // excess away, which produced a valid-looking entry and, on collision, a + // WRONG private key with no error. Refuse loudly instead. + bool overflow = (N.bits64[DIST_WORDS - 1] & (DIST_SIGN_MASK | DIST_TYPE_MASK)) != 0; + for(int i = DIST_WORDS; i < NB64BLOCK; i++) + if(N.bits64[i]) overflow = true; + + if(overflow) { ::printf("\nHashTable::Convert: travelled distance exceeds %d bits.\n" "Interval is too large for the DP entry format (max %d bits).\n" +#if DIST_WORDS == 2 + "Rebuild with -DWIDE_DIST for intervals up to 253 bits.\n" +#endif "Aborting rather than storing a truncated distance.\n", DIST_MAG_BITS,MAX_INTERVAL_BITS); exit(-1); } - D->i64[0] = N.bits64[0]; - D->i64[1] = N.bits64[1]; - D->i64[2] = N.bits64[2]; - D->i64[3] = (N.bits64[3] & DIST_MAG_MASK) | sign | type64; + for(int i = 0; i < DIST_WORDS; i++) + D->i64[i] = N.bits64[i]; + D->i64[DIST_WORDS - 1] = (D->i64[DIST_WORDS - 1] & DIST_MAG_MASK) | sign | type64; *h = (x->bits64[2] & HASH_MASK); @@ -240,7 +245,7 @@ int HashTable::MergeH(uint32_t h,FILE* f1,FILE* f2,FILE* fd,uint32_t* nbDP,uint3 int HashTable::Add(Int *x,Int *d,uint32_t type) { int128_t X; - int256_t D; + dist_t D; uint64_t h; Convert(x,d,type,&h,&X,&D); ENTRY* e = CreateEntry(&X,&D); @@ -258,24 +263,22 @@ void HashTable::ReAllocate(uint64_t h,uint32_t add) { } -int HashTable::Add(uint64_t h,int128_t *x,int256_t *d) { +int HashTable::Add(uint64_t h,int128_t *x,dist_t *d) { ENTRY *e = CreateEntry(x,d); return Add(h,e); } -void HashTable::CalcDistAndType(int256_t d,Int* kDist,uint32_t* kType) { +void HashTable::CalcDistAndType(dist_t d,Int* kDist,uint32_t* kType) { - *kType = (d.i64[3] & DIST_TYPE_MASK) != 0; - int sign = (d.i64[3] & DIST_SIGN_MASK) != 0; - d.i64[3] &= DIST_MAG_MASK; + *kType = (d.i64[DIST_WORDS - 1] & DIST_TYPE_MASK) != 0; + int sign = (d.i64[DIST_WORDS - 1] & DIST_SIGN_MASK) != 0; + d.i64[DIST_WORDS - 1] &= DIST_MAG_MASK; kDist->SetInt32(0); - kDist->bits64[0] = d.i64[0]; - kDist->bits64[1] = d.i64[1]; - kDist->bits64[2] = d.i64[2]; - kDist->bits64[3] = d.i64[3]; + for(int i = 0; i < DIST_WORDS; i++) + kDist->bits64[i] = d.i64[i]; if(sign) kDist->ModNegK1order(); } @@ -403,7 +406,7 @@ void HashTable::SaveTable(FILE* f,uint32_t from,uint32_t to,bool printPoint) { fwrite(&E[h].maxItem,sizeof(uint32_t),1,f); for(uint32_t i = 0; i < E[h].nbItem; i++) { fwrite(&(E[h].items[i]->x),sizeof(int128_t),1,f); - fwrite(&(E[h].items[i]->d),sizeof(int256_t),1,f); + fwrite(&(E[h].items[i]->d),sizeof(dist_t),1,f); if(printPoint) { pointPrint++; if(pointPrint > point) { @@ -473,7 +476,7 @@ void HashTable::LoadTable(FILE* f,uint32_t from,uint32_t to) { for(uint32_t i = 0; i < E[h].nbItem; i++) { ENTRY* e = (ENTRY*)malloc(sizeof(ENTRY)); fread(&(e->x),sizeof(int128_t),1,f); - fread(&(e->d),sizeof(int256_t),1,f); + fread(&(e->d),sizeof(dist_t),1,f); E[h].items[i] = e; } diff --git a/HashTable.h b/HashTable.h index 9c1e4d1e..af30cb11 100644 --- a/HashTable.h +++ b/HashTable.h @@ -20,6 +20,7 @@ #include #include +#include "Constants.h" #include "SECPK1/Point.h" #ifdef WIN64 #include @@ -57,30 +58,36 @@ typedef union int256_s int256_t; #define safe_free(x) if(x) {free(x);x=NULL;} -// Distance field width. The sign and kangaroo-type flags live in the top two -// bits, exactly as in the original 128bit layout, so the magnitude gets -// b253..b0. That is 254 bits for a search that needs at most ~160, i.e. plenty -// of headroom -- moving the flags out to a separate byte would only pad ENTRY -// from 48 to 56 bytes and grow every DP table by 17% for nothing. -#define DIST_MAG_BITS 254 +// Distance field width -- selected by WIDE_DIST in Constants.h. The sign and +// kangaroo-type flags always live in the top two bits of the top word, so +// everything below is written once and works at either width. +#if DIST_WORDS == 4 +typedef int256_t dist_t; #define MAX_INTERVAL_BITS 253 -#define DIST_SIGN_MASK 0x8000000000000000ULL // b255, in i64[3] -#define DIST_TYPE_MASK 0x4000000000000000ULL // b254, in i64[3] -#define DIST_MAG_MASK 0x3FFFFFFFFFFFFFFFULL // b253..b192, in i64[3] +#else +typedef int128_t dist_t; +#define MAX_INTERVAL_BITS 125 +#endif + +#define DIST_MAG_BITS (64 * DIST_WORDS - 2) +#define DIST_SIGN_MASK 0x8000000000000000ULL // top bit of the top word +#define DIST_TYPE_MASK 0x4000000000000000ULL +#define DIST_MAG_MASK 0x3FFFFFFFFFFFFFFFULL // We store only 128 (+18) bit a the x value which give a probabilty a wrong collision after 2^73 entries typedef struct { int128_t x; // Poisition of kangaroo (128bit LSB) - int256_t d; // Travelled distance (b255=sign b254=kangaroo type, b253..b0 distance + dist_t d; // Travelled distance: top bit sign, next bit kangaroo type, + // remaining DIST_MAG_BITS the magnitude } ENTRY; -// On-disk and on-wire size of one ENTRY. Hard-coded rather than sizeof() at -// each use site so a layout change cannot silently reinterpret old files. -#define ENTRY_SIZE 48 -static_assert(sizeof(ENTRY) == ENTRY_SIZE,"ENTRY must stay packed at 48 bytes"); +// On-disk and on-wire size of one ENTRY. Derived, then asserted, so a layout +// change cannot silently reinterpret old files. +#define ENTRY_SIZE (16 + 8 * DIST_WORDS) +static_assert(sizeof(ENTRY) == ENTRY_SIZE,"ENTRY layout has padding"); typedef struct { @@ -96,7 +103,7 @@ class HashTable { HashTable(); int Add(Int *x,Int *d,uint32_t type); - int Add(uint64_t h,int128_t *x,int256_t *d); + int Add(uint64_t h,int128_t *x,dist_t *d); int Add(uint64_t h,ENTRY *e); uint64_t GetNbItem(); void Reset(); @@ -115,16 +122,16 @@ class HashTable { Int kDist; uint32_t kType; - static void Convert(Int *x,Int *d,uint32_t type,uint64_t *h,int128_t *X,int256_t *D); + static void Convert(Int *x,Int *d,uint32_t type,uint64_t *h,int128_t *X,dist_t *D); static int MergeH(uint32_t h,FILE* f1,FILE* f2,FILE* fd,uint32_t *nbDP,uint32_t* duplicate, Int* d1,uint32_t* k1,Int* d2,uint32_t* k2); - static void CalcDistAndType(int256_t d,Int* kDist,uint32_t* kType); + static void CalcDistAndType(dist_t d,Int* kDist,uint32_t* kType); private: - ENTRY *CreateEntry(int128_t *x,int256_t *d); + ENTRY *CreateEntry(int128_t *x,dist_t *d); static int compare(int128_t *i1,int128_t *i2); - static bool sameDist(int256_t *a,int256_t *b); + static bool sameDist(dist_t *a,dist_t *b); std::string GetStr(int128_t *i); }; diff --git a/Kangaroo.cpp b/Kangaroo.cpp index 747c3ba9..ffea7461 100644 --- a/Kangaroo.cpp +++ b/Kangaroo.cpp @@ -313,7 +313,7 @@ bool Kangaroo::AddToTable(Int *pos,Int *dist,uint32_t kType) { } -bool Kangaroo::AddToTable(uint64_t h,int128_t *x,int256_t *d) { +bool Kangaroo::AddToTable(uint64_t h,int128_t *x,dist_t *d) { int addStatus = hashTable.Add(h,x,d); if(addStatus== ADD_COLLISION) { diff --git a/Kangaroo.h b/Kangaroo.h index 6d93b415..c148563e 100644 --- a/Kangaroo.h +++ b/Kangaroo.h @@ -96,7 +96,7 @@ typedef struct { uint32_t kIdx; uint32_t h; int128_t x; - int256_t d; + dist_t d; } DP; @@ -117,12 +117,20 @@ typedef struct { } DP_CACHE; // Work file type -// Bumped for the 256bit distance field: ENTRY grew from 32 to 48 bytes, so a -// pre-existing work file read with the new layout would silently desynchronise -// and produce garbage DPs. Old files are now rejected by magic instead. +// The magics depend on the distance width. A WIDE_DIST build writes 48-byte +// entries; reading those with the 32-byte layout (or vice versa) would +// silently desynchronise and yield garbage DPs, so a mismatched file is +// rejected by magic instead. A default build keeps the original values and +// stays byte-compatible with upstream work files. +#ifdef WIDE_DIST #define HEADW 0xFA6A8011 // Full work file (256bit distance) #define HEADK 0xFA6A8012 // Kangaroo only file (256bit distance) #define HEADKS 0xFA6A8013 // Compressed Kangaroo only file (256bit distance) +#else +#define HEADW 0xFA6A8001 // Full work file +#define HEADK 0xFA6A8002 // Kangaroo only file +#define HEADKS 0xFA6A8003 // Compressed Kangaroo only file +#endif // Number of Hash entry per partition #define H_PER_PART (HASH_SIZE / MERGE_PART) @@ -168,7 +176,7 @@ class Kangaroo { void SetDP(int size); void CreateHerd(int nbKangaroo,Int *px, Int *py, Int *d, int firstType,bool lock=true); void CreateJumpTable(); - bool AddToTable(uint64_t h,int128_t *x,int256_t *d); + bool AddToTable(uint64_t h,int128_t *x,dist_t *d); bool AddToTable(Int *pos,Int *dist,uint32_t kType); bool SendToServer(std::vector &dp,uint32_t threadId,uint32_t gpuId); bool CheckKey(Int d1,Int d2,uint8_t type); @@ -184,7 +192,7 @@ class Kangaroo { void SaveWork(uint64_t totalCount,double totalTime,TH_PARAM *threads,int nbThread); void SaveServerWork(); void FetchWalks(uint64_t nbWalk,Int *x,Int *y,Int *d); - void FetchWalks(uint64_t nbWalk,std::vector& kangs,Int* x,Int* y,Int* d); + void FetchWalks(uint64_t nbWalk,std::vector& kangs,Int* x,Int* y,Int* d); void FectchKangaroos(TH_PARAM *threads); FILE *ReadHeader(std::string fileName,uint32_t *version,int type); bool SaveHeader(std::string fileName,FILE* f,int type,uint64_t totalCount,double totalTime); @@ -207,8 +215,8 @@ class Kangaroo { void InitSocket(); void WaitForServer(); int32_t GetServerStatus(); - bool SendKangaroosToServer(std::string& fileName,std::vector& kangs); - bool GetKangaroosFromServer(std::string& fileName,std::vector& kangs); + bool SendKangaroosToServer(std::string& fileName,std::vector& kangs); + bool GetKangaroosFromServer(std::string& fileName,std::vector& kangs); #ifdef WIN64 HANDLE ghMutex; diff --git a/Network.cpp b/Network.cpp index baf34bce..b16b2027 100644 --- a/Network.cpp +++ b/Network.cpp @@ -44,9 +44,14 @@ static SOCKET serverSock = 0; #define SERVER_VERSION 3 -// Bumped with the 256bit distance: DP grew 40 -> 56 bytes and the kangaroo -// block grew 16 -> 32 bytes per kangaroo. An old peer would misparse both. +// The protocol magic depends on the distance width: a WIDE_DIST build sends +// 56-byte DPs and 32-byte kangaroo blocks, a default build 40 and 16. A +// mismatched peer must fail the handshake rather than misparse every packet. +#ifdef WIDE_DIST #define SERVER_HEADER 0x67DEDDD1 +#else +#define SERVER_HEADER 0x67DEDDC1 +#endif #define KANG_PER_BLOCK 2048 @@ -338,7 +343,7 @@ bool Kangaroo::HandleRequest(TH_PARAM *p) { uint64_t nbKangaroo = 0; uint32_t strSize; char fileName[256]; - int256_t* KBuff; + dist_t* KBuff; uint32_t nbK; uint32_t header = HEADKS; uint32_t version = 0; @@ -380,7 +385,7 @@ bool Kangaroo::HandleRequest(TH_PARAM *p) { PUT("nbKangaroo",p->clientSock,&nbKangaroo,sizeof(uint64_t),ntimeout); checkSum.SetInt32(0); - KBuff = (int256_t*)malloc(KANG_PER_BLOCK * sizeof(int256_t)); + KBuff = (dist_t*)malloc(KANG_PER_BLOCK * sizeof(dist_t)); while(nbKangaroo > 0) { @@ -391,17 +396,15 @@ bool Kangaroo::HandleRequest(TH_PARAM *p) { } for(uint32_t k = 0; k < nbK; k++) { - ::fread(&KBuff[k],32,1,f); + ::fread(&KBuff[k],sizeof(dist_t),1,f); // Checksum K.SetInt32(0); - K.bits64[1] = KBuff[k].i64[1]; - K.bits64[0] = KBuff[k].i64[0]; - K.bits64[2] = KBuff[k].i64[2]; - K.bits64[3] = KBuff[k].i64[3]; + for(int w = 0; w < DIST_WORDS; w++) + K.bits64[w] = KBuff[k].i64[w]; checkSum.Add(&K); } - PUTFREE("packet",p->clientSock,KBuff,nbK * 32,ntimeout,KBuff); + PUTFREE("packet",p->clientSock,KBuff,nbK * (8*DIST_WORDS),ntimeout,KBuff); nbKangaroo -= nbK; @@ -429,7 +432,7 @@ bool Kangaroo::HandleRequest(TH_PARAM *p) { uint32_t fileNameSize; char fileNameTmp[264]; char fileName[256]; - int256_t *KBuff; + dist_t *KBuff; uint32_t nbK; uint32_t header = HEADKS; uint32_t version = 0; @@ -466,7 +469,7 @@ bool Kangaroo::HandleRequest(TH_PARAM *p) { checkSum.SetInt32(0); - KBuff = (int256_t *)malloc(KANG_PER_BLOCK*sizeof(int256_t)); + KBuff = (dist_t *)malloc(KANG_PER_BLOCK*sizeof(dist_t)); while(nbKangaroo>0) { @@ -476,16 +479,14 @@ bool Kangaroo::HandleRequest(TH_PARAM *p) { nbK = (uint32_t)nbKangaroo; } - GETFREE("packet",p->clientSock,KBuff,nbK * 32,ntimeout,KBuff); + GETFREE("packet",p->clientSock,KBuff,nbK * (8*DIST_WORDS),ntimeout,KBuff); for(uint32_t k = 0; k < nbK; k++) { - ::fwrite(&KBuff[k],32,1,f); + ::fwrite(&KBuff[k],sizeof(dist_t),1,f); // Checksum K.SetInt32(0); - K.bits64[1] = KBuff[k].i64[1]; - K.bits64[0] = KBuff[k].i64[0]; - K.bits64[2] = KBuff[k].i64[2]; - K.bits64[3] = KBuff[k].i64[3]; + for(int w = 0; w < DIST_WORDS; w++) + K.bits64[w] = KBuff[k].i64[w]; checkSum.Add(&K); } @@ -708,7 +709,7 @@ void Kangaroo::RunServer() { } SetDP(initDPSize); - if(sizeof(DP) != 56) { + if(sizeof(DP) != (int)(8 + 16 + (8*DIST_WORDS))) { ::printf("Error: Invalid DP size struct\n"); exit(-1); } @@ -992,14 +993,14 @@ void Kangaroo::WaitForServer() { } // Get Kangaroo from server -bool Kangaroo::GetKangaroosFromServer(std::string& fileName,std::vector& kangs) { +bool Kangaroo::GetKangaroosFromServer(std::string& fileName,std::vector& kangs) { int nbRead; int nbWrite; uint32_t fileNameSize = (uint32_t)fileName.length(); uint64_t nbKangaroo = 0; uint32_t nbK; - int256_t* KBuff; + dist_t* KBuff; Int checkSum; WaitForServer(); @@ -1023,7 +1024,7 @@ bool Kangaroo::GetKangaroosFromServer(std::string& fileName,std::vector& kangs) { +bool Kangaroo::SendKangaroosToServer(std::string& fileName,std::vector& kangs) { int nbWrite; uint32_t fileNameSize = (uint32_t)fileName.length(); uint64_t nbKangaroo = kangs.size(); uint64_t pos; uint32_t nbK; - int256_t *KBuff; + dist_t *KBuff; Int checkSum; WaitForServer(); @@ -1105,7 +1104,7 @@ bool Kangaroo::SendKangaroosToServer(std::string& fileName,std::vector PUT("fileName",serverConn,fileName.c_str(),fileNameSize,ntimeout); PUT("nbKangaroo",serverConn,&nbKangaroo,sizeof(uint64_t),ntimeout); - KBuff = (int256_t*)malloc(KANG_PER_BLOCK * sizeof(int256_t)); + KBuff = (dist_t*)malloc(KANG_PER_BLOCK * sizeof(dist_t)); checkSum.SetInt32(0); pos = 0; @@ -1124,19 +1123,17 @@ bool Kangaroo::SendKangaroosToServer(std::string& fileName,std::vector } for(uint32_t k = 0; k < nbK; k++) { - memcpy(&KBuff[k],&kangs[pos],32); + memcpy(&KBuff[k],&kangs[pos],sizeof(dist_t)); pos++; // Checksum Int K; K.SetInt32(0); - K.bits64[1] = KBuff[k].i64[1]; - K.bits64[0] = KBuff[k].i64[0]; - K.bits64[2] = KBuff[k].i64[2]; - K.bits64[3] = KBuff[k].i64[3]; + for(int w = 0; w < DIST_WORDS; w++) + K.bits64[w] = KBuff[k].i64[w]; checkSum.Add(&K); } - PUTFREE("packet",serverConn,KBuff,nbK * 32,ntimeout,KBuff); + PUTFREE("packet",serverConn,KBuff,nbK * (8*DIST_WORDS),ntimeout,KBuff); nbKangaroo -= nbK; @@ -1176,7 +1173,7 @@ bool Kangaroo::SendToServer(std::vector &dps,uint32_t threadId,uint32_t gp for(uint32_t i = 0; i &dps,uint32_t threadId,uint32_t gp dp[i].h = (uint32_t)h; dp[i].x.i64[0] = X.i64[0]; dp[i].x.i64[1] = X.i64[1]; - dp[i].d.i64[0] = D.i64[0]; - dp[i].d.i64[1] = D.i64[1]; - dp[i].d.i64[2] = D.i64[2]; - dp[i].d.i64[3] = D.i64[3]; + for(int w = 0; w < DIST_WORDS; w++) + dp[i].d.i64[w] = D.i64[w]; } diff --git a/README.md b/README.md index ad8f8cb0..6fc73a94 100644 --- a/README.md +++ b/README.md @@ -1,7 +1,13 @@ # Pollard's kangaroo for SECPK1 A Pollard's kangaroo interval ECDLP solver for SECP256K1 (based on VanitySearch engine).\ -**This program is limited to a 125bit interval search.** +**By default this program is limited to a 125bit interval search.** Build with +`-DWIDE_DIST` (or uncomment `WIDE_DIST` in `Constants.h`) to raise that to 253 +bits, which is what puzzle 130 and above need. See +[Wider intervals](#wider-intervals-wide_dist) below for the trade-offs. + +Exceeding the compiled limit is refused with an error. It used to be accepted +silently and produce a wrong private key. # Feature @@ -74,6 +80,53 @@ ex 0335BB25364370D4DD14A9FC2B406D398C4B53C85BE58FCC7297BD34004602EBEC ``` +# Wider intervals (WIDE_DIST) + +The travelled distance is stored alongside each distinguished point. By default +that field is 128 bits: the top bit is the sign, the next the kangaroo type, +leaving 126 bits of magnitude, which caps the search interval at 125 bits. + +Every unsolved kangaroo-able puzzle is above that cap -- #135 is a 2^134 +interval, #140 is 2^139. To search them, build with `-DWIDE_DIST` (or uncomment +`WIDE_DIST` in `Constants.h`) for a 256-bit field: 254 bits of magnitude, +intervals up to 253 bits. + +| | default | `-DWIDE_DIST` | +|------------------------|-----------|---------------| +| max interval | 125 bits | 253 bits | +| DP entry | 32 bytes | 48 bytes | +| kangaroo (device) | 10 words | 12 words | +| DP network packet | 40 bytes | 56 bytes | +| kangaroo transfer | 16 bytes | 32 bytes | + +So a wide build costs about 50% more RAM for the same number of DPs, 20% more +device memory, and measured 4-7% throughput on an RTX 3050 Laptop. + +Work files and the client/server protocol are **not** interchangeable between +the two builds. The format magics differ, so a mismatched work file is rejected +with "Not a work file" and a mismatched peer fails the handshake -- neither is +silently misparsed. A default build reads and writes exactly the same format as +upstream. + +If you exceed the compiled interval limit the program now stops: + +``` +HashTable::Convert: travelled distance exceeds 126 bits. +Interval is too large for the DP entry format (max 125 bits). +Rebuild with -DWIDE_DIST for intervals up to 253 bits. +Aborting rather than storing a truncated distance. +``` + +Previously the excess bits were masked off, the entry looked valid, and a +collision produced a wrong private key with no warning of any kind. + +Two scripts check the layout without needing a compiler or a GPU: + +``` +python tools/check_layout.py # producer/consumer offsets, both widths +python tools/test_distfield.py # bit packing round-trip, both widths +``` + # Note on Time/Memory tradeoff of the DP method The distinguished point (DP) method is an efficient method for storing random walks and detect collision between them. Instead of storing all points of all kangagroo's random walks, we store only points that have an x value starting with dpBit zero bits. When 2 kangaroos collide, they will then follow the same path because their jumps are a function of their x values. The collision will be then detected when the 2 kangaroos reach a distinguished point.\ diff --git a/build_gpu.bat b/build_gpu.bat index 210d7d12..6a6cff67 100644 --- a/build_gpu.bat +++ b/build_gpu.bat @@ -6,6 +6,16 @@ rem sm_86 = GA107 (RTX 3050 Laptop). Override with: build_gpu.bat 89 set CCAP=%1 if "%CCAP%"=="" set CCAP=86 +rem second arg "wide" builds the 256-bit distance variant +set WIDE= +if /i "%2"=="wide" set WIDE=-DWIDE_DIST +set WIDEC= +if /i "%2"=="wide" set WIDEC=/DWIDE_DIST +set OUTEXE=kangaroo-gpu.exe +if /i "%2"=="wide" set OUTEXE=kangaroo-wide.exe +set OBJD=objgpu +if /i "%2"=="wide" set OBJD=objgpuw + rem CUDA 13.x may reject the newest MSVC as an unsupported host compiler. rem MSVC_DIR pins nvcc to a known-good toolset; leave empty to use the default. set "MSVC_DIR=C:\Program Files (x86)\Microsoft Visual Studio\18\BuildTools\VC\Tools\MSVC\14.44.35207\bin\Hostx64\x64" @@ -27,27 +37,27 @@ if not exist "%CUDA_PATH%\bin\nvcc.exe" (echo nvcc NOT FOUND under %CUDA_PATH% & echo Using "%CUDA_PATH%", sm_%CCAP% rem separate obj dir: the CPU build left non-WITHGPU objects in obj\ -if not exist objgpu mkdir objgpu -del /q objgpu\*.obj 2>nul +if not exist %OBJD% mkdir %OBJD% +del /q %OBJD%\*.obj 2>nul set CCBIN= if exist "%MSVC_DIR%\cl.exe" set CCBIN=-ccbin "%MSVC_DIR%" "%CUDA_PATH%\bin\nvcc.exe" -maxrregcount=0 --ptxas-options=-v --compile %CCBIN% ^ - -m64 -O2 -I. -I"%CUDA_PATH%\include" -DWITHGPU -DWIN64 -D_CRT_SECURE_NO_WARNINGS ^ + -m64 -O2 -I. -I"%CUDA_PATH%\include" -DWITHGPU -DWIN64 %WIDE% -D_CRT_SECURE_NO_WARNINGS ^ -gencode=arch=compute_%CCAP%,code=sm_%CCAP% ^ - -o objgpu\GPUEngine.obj -c GPU\GPUEngine.cu + -o %OBJD%\GPUEngine.obj -c GPU\GPUEngine.cu if errorlevel 1 (echo NVCC FAILED & exit /b 1) -cl /nologo /EHsc /O2 /std:c++14 /DWITHGPU /DWIN64 /D_CRT_SECURE_NO_WARNINGS ^ +cl /nologo /EHsc /O2 /std:c++14 /DWITHGPU /DWIN64 %WIDEC% /D_CRT_SECURE_NO_WARNINGS ^ /I. /I"%CUDA_PATH%\include" /c ^ main.cpp Kangaroo.cpp HashTable.cpp Backup.cpp Thread.cpp Check.cpp ^ Network.cpp Merge.cpp PartMerge.cpp Timer.cpp ^ SECPK1\Int.cpp SECPK1\IntGroup.cpp SECPK1\IntMod.cpp ^ - SECPK1\Point.cpp SECPK1\Random.cpp SECPK1\SECP256K1.cpp /Foobjgpu\ + SECPK1\Point.cpp SECPK1\Random.cpp SECPK1\SECP256K1.cpp /Fo%OBJD%\ if errorlevel 1 (echo COMPILE FAILED & exit /b 1) -link /nologo /OUT:kangaroo-gpu.exe objgpu\*.obj ws2_32.lib advapi32.lib ^ +link /nologo /OUT:%OUTEXE% %OBJD%\*.obj ws2_32.lib advapi32.lib ^ /LIBPATH:"%CUDA_PATH%\lib\x64" cudart.lib if errorlevel 1 (echo LINK FAILED & exit /b 1) echo BUILD OK diff --git a/tools/check_layout.py b/tools/check_layout.py index 4f78e151..7bc8e3b6 100644 --- a/tools/check_layout.py +++ b/tools/check_layout.py @@ -1,162 +1,171 @@ #!/usr/bin/env python3 -"""Static consistency check for the 256-bit distance patch. - -Phase 2 is mostly hand-edited offsets across the CUDA kernel, the host readout -and two wire formats. A compiler catches type errors; it does NOT catch a DP -written at word 9 and read at word 8. This parses the actual sources and -asserts producer and consumer agree. +"""Static consistency check for the configurable distance-width patch. + +The width is a compile-time switch (DIST_WORDS = 2 by default, 4 with +WIDE_DIST), and most of the affected code is hand-written offsets across the +CUDA kernel, the host readout and two wire formats. A compiler catches type +errors; it does NOT catch a DP written at word 9 and read at word 8, and it +only ever checks the width you happened to build. This parses the sources and +asserts producer and consumer agree at BOTH widths. """ import re -import sys from pathlib import Path -# Works both from inside the repo (tools/) and from a parent dir holding -# a "Kangaroo" checkout. _here = Path(__file__).resolve().parent ROOT = _here.parent if (_here.parent / "HashTable.h").exists() else _here.parent / "Kangaroo" +WIDTHS = (2, 4) + def read(rel): return (ROOT / rel).read_text(encoding="utf-8", errors="replace") -def define(src, name): - m = re.search(r"^#define\s+%s\s+(\d+)" % re.escape(name), src, re.M) - assert m, "no #define %s" % name - return int(m.group(1)) +def evaluate(expr, dw): + """Evaluate a C constant expression that may mention DIST_WORDS.""" + expr = expr.replace("DIST_WORDS", str(dw)) + assert re.fullmatch(r"[\d\s()+*/-]+", expr), "unexpected tokens: %r" % expr + return eval(expr) -def check_item_layout(): +def define_expr(src, name): + m = re.search(r"^#define\s+%s\s+(.+?)\s*(?://.*)?$" % re.escape(name), src, re.M) + assert m, "no #define %s" % name + return m.group(1).strip() + + +def active_block(src, dw): + """Strip the inactive arm of every `#if DIST_WORDS == 4` / #else block.""" + out, state = [], [] + for line in src.split("\n"): + st = line.strip() + if st.startswith("#if DIST_WORDS == 4"): + state.append(dw == 4) + continue + if state and st == "#else": + state[-1] = not state[-1] + continue + if state and st == "#endif": + state.pop() + continue + if all(state): + out.append(line) + return "\n".join(out) + + +def check_item_layout(dw): """GPU DP output: OutputDP writer vs GPUEngine.cu reader.""" eng_h = read("GPU/GPUEngine.h") - item_size = define(eng_h, "ITEM_SIZE") - item32 = item_size // 4 - - # x[4 u64] + d[4 u64] + kIdx[1 u64] - assert item_size == (4 + 4 + 1) * 8 == 72, item_size - - math_h = read("GPU/GPUMath.h") + item32 = evaluate(define_expr(eng_h, "ITEM_SIZE32"), dw) + item_size = evaluate(define_expr(eng_h, "ITEM_SIZE").replace("ITEM_SIZE32", str(item32)), dw) + + # x[4 u64] + d[dw u64] + kIdx[1 u64] + assert item_size == (4 + dw + 1) * 8, (dw, item_size) + assert item32 == item_size // 4 + + math_h = active_block(read("GPU/GPUMath.h"), dw) + # inline DP_EXTRA_DIST_WORDS so the writer reads as one flat list + m = re.search(r"#define DP_EXTRA_DIST_WORDS\(d\)(.*?)\n#(?:else|endif)", + read("GPU/GPUMath.h"), re.S) + extra = m.group(1) if (m and dw == 4) else "" body = math_h[math_h.index("#define OutputDP"):] - body = body[:body.index("\n}")] - slots = [(int(o), field, int(i)) - for o, field, i in re.findall( - r"out\[pos\*ITEM_SIZE32 \+ (\d+)\] = \(\(uint32_t \*\)(\w+)\)\[(\d+)\];", - body)] - assert slots, "OutputDP body not parsed" - - # Every uint32 of the item is written exactly once, contiguous from 1. - written = [s[0] for s in slots] - assert written == list(range(1, item32 + 1)), \ - "OutputDP writes %s, expected 1..%d" % (written, item32) - - # Each source field is written low word first, no gaps. - for field, count in (("x", 8), ("d", 8), ("idx", 2)): - idxs = [i for _, f, i in slots if f == field] - assert idxs == list(range(count)), "%s writes %s" % (field, idxs) - - # Reader side: itemPtr = out + i*ITEM_SIZE32 + 1, so writer slot N -> itemPtr[N-1]. - eng_cu = read("GPU/GPUEngine.cu") - x_off = [int(n) for n in re.findall(r"it\.x\.bits64\[(\d)\] = x\[\d\];", eng_cu)] - d_base = int(re.search(r"uint64_t \*d = \(uint64_t \*\)\(itemPtr \+ (\d+)\);", - eng_cu).group(1)) - d_words = re.findall(r"it\.d\.bits64\[(\d)\] = d\[(\d)\];", eng_cu) - k_off = int(re.search(r"it\.kIdx = \*\(\(uint64_t\*\)\(itemPtr \+ (\d+)\)\);", - eng_cu).group(1)) - - assert x_off == [0, 1, 2, 3], x_off - assert d_base == 8, "d read at u32 offset %d, writer put it at 8" % d_base - assert [(int(a), int(b)) for a, b in d_words] == [(0, 0), (1, 1), (2, 2), (3, 3)], \ - "distance words not read 1:1: %s" % d_words - assert k_off == 16, "kIdx read at u32 %d, writer put it at 16" % k_off - # kIdx is the last 2 u32 of the item - assert k_off + 2 == item32, "kIdx at %d + 2 != item32 %d" % (k_off, item32) + body = body[:body.index("\n}")].replace("DP_EXTRA_DIST_WORDS(d)", extra) + + slots = [] + for off, field, idx in re.findall( + r"out\[pos\*ITEM_SIZE32 \+ ([^\]]+)\] = \(\(uint32_t \*\)(\w+)\)\[(\d+)\];", body): + slots.append((evaluate(off, dw), field, int(idx))) + + # every uint32 of the item written exactly once; slot 0 is the found-counter + written = sorted(s[0] for s in slots) + assert len(written) == len(set(written)), (dw, "duplicate slot write", written) + assert written == list(range(1, item32 + 1)), (dw, written, item32) + + # fields land contiguously and in order + for field, base, count in (("x", 1, 8), ("d", 9, 2 * dw), ("idx", 9 + 2 * dw, 2)): + got = sorted((o, i) for o, f, i in slots if f == field) + assert got == [(base + i, i) for i in range(count)], (dw, field, got) + + # reader side must use the same offsets + cu = read("GPU/GPUEngine.cu") + assert "uint64_t *x = (uint64_t *)itemPtr;" in cu + assert "uint64_t *d = (uint64_t *)(itemPtr + 8);" in cu, "d must be read at +8" + m = re.search(r"it\.kIdx = \*\(\(uint64_t\*\)\(itemPtr \+ ([^)]+)\)\);", cu) + assert m and evaluate(m.group(1), dw) == 8 + 2 * dw, (dw, m and m.group(1)) return item_size -def check_kangaroo_slots(): - """Device kangaroo record: every slot used must fit inside KSIZE.""" +def check_kangaroo_slots(dw): + """Device kangaroo record: px[4] py[4] dist[dw] (+lastJump) must fit KSIZE.""" eng_h = read("GPU/GPUEngine.h") - ksize_sym = int(re.search(r"#ifdef USE_SYMMETRY\s*\n#define KSIZE (\d+)", eng_h).group(1)) - ksize = int(re.search(r"#else\s*\n#define KSIZE (\d+)", eng_h).group(1)) - assert (ksize, ksize_sym) == (12, 13), (ksize, ksize_sym) - - math_h = read("GPU/GPUMath.h") - used = {int(n) for n in re.findall(r"IDX \+ (\d+) \* blockDim\.x \+ stride", math_h)} - assert used == set(range(13)), "device slots used: %s" % sorted(used) - assert max(used) < ksize_sym - - # Non-symmetry build must not touch slot 12 (that is the lastJump word). - non_sym = re.sub(r"#ifdef USE_SYMMETRY.*?#endif", "", math_h, flags=re.S) - used_ns = {int(n) for n in re.findall(r"IDX \+ (\d+) \* blockDim\.x \+ stride", non_sym)} - assert max(used_ns) < ksize, "non-symmetry uses slot %d, KSIZE=%d" % (max(used_ns), ksize) - - # Host side writes the same slots. - eng_cu = read("GPU/GPUEngine.cu") - host = {int(n) for n in re.findall(r"t \+ (\d+) \* nbThreadPerGroup", eng_cu)} - assert host == set(range(13)), "host slots: %s" % sorted(host) - - # dist is 4 words everywhere it is declared. - assert "dist[GPU_GRP_SIZE][2]" not in math_h - assert "dist[GPU_GRP_SIZE][2]" not in read("GPU/GPUCompute.h") - # and accumulated with the 4-word carry chain - assert "Add256(dist[g],jD[jmp]);" in read("GPU/GPUCompute.h") - assert "jD[NB_JUMP][4]" in math_h + ks = re.findall(r"#define KSIZE\s+(.+)", eng_h) + assert len(ks) == 2, ks + ksize_sym, ksize = (evaluate(k.strip(), dw) for k in ks) + assert ksize == 8 + dw, (dw, ksize) + assert ksize_sym == ksize + 1, (dw, ksize_sym) + + math_h = active_block(read("GPU/GPUMath.h"), dw) + used = set() + for expr in re.findall(r"\(a\)\[IDX \+ ([^*]+)\* blockDim\.x \+ stride\]", math_h): + used.add(evaluate(expr.strip(), dw)) + assert used, "no kangaroo slot accesses found" + assert max(used) < ksize_sym, (dw, sorted(used), ksize_sym) + assert used >= set(range(0, 8 + dw)), (dw, sorted(used)) + + cu = active_block(read("GPU/GPUEngine.cu"), dw) + host = set() + for expr in re.findall(r"t \+ ([^*]+)\* nbThreadPerGroup", cu): + host.add(evaluate(expr.strip(), dw)) + assert max(host) < ksize_sym, (dw, sorted(host), ksize_sym) return ksize -def check_add256(): - """Carry must chain through all four words: add.cc / addc.cc / addc.cc / addc.""" +def check_add_carry(dw): + """AddDist must select a carry chain covering exactly DIST_WORDS words.""" math_h = read("GPU/GPUMath.h") - body = math_h[math_h.index("#define Add256"):] - body = body[:body.index("(a)[3]);}") + len("(a)[3]);}")] - ops = re.findall(r"(UADDO1|UADDC1|UADD1)\(\(r\)\[(\d)\], \(a\)\[(\d)\]\)", body) - assert [o[0] for o in ops] == ["UADDO1", "UADDC1", "UADDC1", "UADD1"], ops - assert [o[1] for o in ops] == list("0123") and [o[2] for o in ops] == list("0123") + sel = re.search(r"#if DIST_WORDS == 4\s*\n#define AddDist\(r,a\) (\w+)\(r,a\)\s*\n" + r"#else\s*\n#define AddDist\(r,a\) (\w+)\(r,a\)", math_h) + assert sel, "AddDist selector missing" + name = sel.group(1) if dw == 4 else sel.group(2) + body = math_h[math_h.index("#define %s(r,a)" % name):] + body = body[:body.index("}")] + ops = re.findall(r"(UADDO1|UADDC1|UADD1)\(\(r\)\[(\d+)\], \(a\)\[(\d+)\]\)", body) + assert len(ops) == dw, (dw, name, ops) + assert [int(o[1]) for o in ops] == list(range(dw)) + assert [int(o[2]) for o in ops] == list(range(dw)) + # carry must start, chain, then terminate + assert ops[0][0] == "UADDO1" and ops[-1][0] == "UADD1", ops + assert all(o[0] == "UADDC1" for o in ops[1:-1]), ops + + comp = read("GPU/GPUCompute.h") + assert "AddDist(dist[g],jD[jmp]);" in comp + assert "uint64_t dist[GPU_GRP_SIZE][DIST_WORDS];" in comp + assert "jD[NB_JUMP][DIST_WORDS]" in math_h + + +def check_host_sizes(dw): + ht = read("HashTable.h") + entry = evaluate(define_expr(ht, "ENTRY_SIZE"), dw) + assert entry == 16 + 8 * dw, (dw, entry) + assert evaluate(define_expr(ht, "DIST_MAG_BITS"), dw) == 64 * dw - 2 + net = read("Network.cpp") + m = re.search(r"if\(sizeof\(DP\) != \(int\)\((.+?)\)\) \{", net) + assert m, "DP size assertion missing" + dp_size = evaluate(m.group(1), dw) + assert dp_size == 8 + 16 + 8 * dw, (dw, dp_size) -def check_wire(): - """DP packet and kangaroo block sizes agree with their runtime assertions.""" - kh = read("Kangaroo.h") - dp = kh[kh.index("// DP transfered over the network"):] - dp = dp[:dp.index("} DP;")] - assert "int128_t x;" in dp and "int256_t d;" in dp, dp - dp_size = 4 + 4 + 16 + 32 # kIdx + h + x + d - assert dp_size == 56 + # kangaroo blocks sized from the type, not a literal + assert net.count("nbK * (8*DIST_WORDS),ntimeout") == 4 + assert "memcpy(&KBuff[k],&kangs[pos],sizeof(dist_t));" in net + assert "kangs.size()*(8*DIST_WORDS) + 16" in read("Backup.cpp") + return entry, dp_size + +def check_checksum_truncation(): + """Only 32 bytes go on the wire; the flags push a herd past 2^256.""" net = read("Network.cpp") - asserted = int(re.search(r"if\(sizeof\(DP\) != (\d+)\)", net).group(1)) - assert asserted == dp_size, "runtime check says %d, struct is %d" % (asserted, dp_size) - - # All four distance words are packed into the outgoing DP. - packed = re.findall(r"dp\[i\]\.d\.i64\[(\d)\] = D\.i64\[(\d)\];", net) - assert [(int(a), int(b)) for a, b in packed] == [(0, 0), (1, 1), (2, 2), (3, 3)], packed - - # Kangaroo block: 32 bytes per kangaroo on every path, none left at 16. - assert "int128_t* KBuff" not in net and "int128_t *KBuff" not in net - # Declaration, cast and allocation size must all agree -- a stale cast here - # is a type error the eye slides right over. - allocs = re.findall(r"KBuff = \((\w+) ?\*\)malloc\(KANG_PER_BLOCK ?\* ?sizeof\((\w+)\)\)", net) - assert len(allocs) == 4, allocs - assert all(c == "int256_t" and t == "int256_t" for c, t in allocs), allocs - for pat in (r"::fread\(&KBuff\[k\],(\d+),1,f\);", - r"::fwrite\(&KBuff\[k\],(\d+),1,f\);", - r"memcpy\(&KBuff\[k\],&kangs\[pos\],(\d+)\);"): - sizes = re.findall(pat, net) - assert sizes and all(s == "32" for s in sizes), (pat, sizes) - pkt = re.findall(r"KBuff,nbK \* (\d+),ntimeout", net) - assert pkt and all(s == "32" for s in pkt), pkt - - # Checksum must cover all four words wherever it is computed. - blocks = re.findall(r"K\.SetInt32\(0\);(.*?)checkSum\.Add\(&K\);", net, re.S) - assert blocks, "no checksum blocks found" - for b in blocks: - got = sorted(int(n) for n in re.findall(r"K\.bits64\[(\d)\] = KBuff", b)) - assert got == [0, 1, 2, 3], "checksum covers words %s" % got - - # The checksum accumulator is a 5-word Int but only 32 bytes go on the wire, - # and the distance now carries flags at b254/b255, so a herd overflows past - # 256 bits. Both sides must truncate or every transfer fails its checksum. sends = len(re.findall(r'PUT\("check[Ss]um",\w+(?:->clientSock)?,checkSum\.bits64,32', net)) compares = len(re.findall(r"if\(!K\.IsEqual\(&checkSum\)\)", net)) truncs = len(re.findall(r"checkSum\.bits64\[4\] = 0;", net)) @@ -165,61 +174,55 @@ def check_wire(): for m in re.finditer(r"checkSum\.bits64\[4\] = 0;(.{0,400})", net, re.S): assert ("IsEqual(&checkSum)" in m.group(1) or "checkSum.bits64,32" in m.group(1)), "truncation not before use" - - assert "kangs.size()*32 + 16" in read("Backup.cpp") - return dp_size, len(blocks) + loops = re.findall(r"for\(int w = 0; w < DIST_WORDS; w\+\+\)\s*\n\s*" + r"K\.bits64\[w\] = KBuff\[k\]\.i64\[w\];", net) + assert len(loops) == 4, len(loops) -def check_entry(): - ht = read("HashTable.h") - assert define(ht, "ENTRY_SIZE") == 16 + 32 == 48 - assert define(ht, "MAX_INTERVAL_BITS") == 253 - assert "int256_t d;" in ht - # Every ENTRY-sized transfer goes through ENTRY_SIZE, and the two halves of - # an ENTRY are read/written at their own widths. Bare 32s elsewhere in these - # files are 256-bit Int header fields (range start/end, key x/y) -- correct - # as-is, so this checks the ENTRY sites specifically rather than banning 32. - htc = read("HashTable.cpp") - assert "::fread(items+i,ENTRY_SIZE,1,f);" in read("Check.cpp") - for pat in (r"::fread\(&e1,ENTRY_SIZE,1,f1\)", r"::fread\(&e2,ENTRY_SIZE,1,f2\)", - r"::fwrite\(output,ENTRY_SIZE,nbd,fd\)", - r"uint64_t hSize = \(uint64_t\)ENTRY_SIZE \* E\[h\]\.nbItem;"): - assert re.search(pat, htc), pat - assert len(re.findall(r"memcpy\(output ?\+ ?nbd,&e\d,ENTRY_SIZE\)", htc)) == 5 - for half, width in (("x", "int128_t"), ("d", "int256_t")): - assert ("fwrite(&(E[h].items[i]->%s),sizeof(%s),1,f)" % (half, width)) in htc - assert ("fread(&(e->%s),sizeof(%s),1,f)" % (half, width)) in htc - # The truncating mask is gone; the guard replaced it. - htc = read("HashTable.cpp") - assert "exit(-1);" in htc and "0xC000000000000000ULL" in htc - - -def check_format_magic(): - """Old files and old peers must be rejected, not misparsed.""" +def check_format_magics(): + """A width mismatch must be refused by magic, never misparsed. The default + build must keep upstream's values so its work files stay compatible.""" kh = read("Kangaroo.h") - for name, old in (("HEADW", "0xFA6A8001"), ("HEADK", "0xFA6A8002"), - ("HEADKS", "0xFA6A8003")): - m = re.search(r"#define %s\s+(0x[0-9A-Fa-f]+)" % name, kh) - assert m and m.group(1).lower() != old.lower(), "%s not bumped" % name - net = read("Network.cpp") - m = re.search(r"#define SERVER_HEADER (0x[0-9A-Fa-f]+)", net) - assert m and m.group(1).lower() != "0x67deddc1", "SERVER_HEADER not bumped" + m = re.search(r"#ifdef WIDE_DIST(.*?)#else(.*?)#endif", kh, re.S) + assert m, "work-file magics are not width-dependent" + wide, narrow = m.group(1), m.group(2) + for name, upstream in (("HEADW", "0xfa6a8001"), ("HEADK", "0xfa6a8002"), + ("HEADKS", "0xfa6a8003")): + w = re.search(r"#define %s\s+(0x[0-9A-Fa-f]+)" % name, wide).group(1) + n = re.search(r"#define %s\s+(0x[0-9A-Fa-f]+)" % name, narrow).group(1) + assert n.lower() == upstream, (name, n, "default must match upstream") + assert w.lower() != n.lower(), (name, "wide magic must differ") - -def main(): - item = check_item_layout() - ksize = check_kangaroo_slots() - check_add256() - dp_size, nblocks = check_wire() - check_entry() - check_format_magic() - print("layout self-check OK") - print(" ENTRY 48 bytes (was 32) x:16 d:32") - print(" ITEM %d bytes (was 56) x:32 d:32 kIdx:8" % item) - print(" KSIZE %d words (was 10) px:4 py:4 dist:4" % ksize) - print(" DP packet %d bytes (was 40)" % dp_size) - print(" kangaroo 32 bytes (was 16), %d checksum sites widened" % nblocks) + net = read("Network.cpp") + m = re.search(r"#ifdef WIDE_DIST\s*\n#define SERVER_HEADER (0x\w+)\s*\n#else\s*\n" + r"#define SERVER_HEADER (0x\w+)", net) + assert m, "protocol magic is not width-dependent" + assert m.group(1) != m.group(2) + assert m.group(2).lower() == "0x67deddc1", (m.group(2), "default must match upstream") + + +def demo(): + results = {} + for dw in WIDTHS: + item = check_item_layout(dw) + ksize = check_kangaroo_slots(dw) + check_add_carry(dw) + entry, dp = check_host_sizes(dw) + results[dw] = (entry, item, ksize, dp, 8 * dw) + check_checksum_truncation() + check_format_magics() + + # the default width must reproduce upstream's layout exactly + assert results[2] == (32, 56, 10, 40, 16), results[2] + assert results[4] == (48, 72, 12, 56, 32), results[4] + + print("layout self-check OK (both widths)") + print(" %-11s %8s %10s" % ("", "default", "WIDE_DIST")) + print(" %-11s %8d %10d" % ("DIST_WORDS", 2, 4)) + for i, name in enumerate(("ENTRY", "ITEM", "KSIZE", "DP packet", "kangaroo")): + print(" %-11s %8d %10d" % (name, results[2][i], results[4][i])) + print(" default reproduces upstream exactly (ENTRY 32, ITEM 56, KSIZE 10, DP 40)") if __name__ == "__main__": - sys.exit(main()) + demo() diff --git a/tools/test_distfield.py b/tools/test_distfield.py index fa150b1d..bc33d09a 100644 --- a/tools/test_distfield.py +++ b/tools/test_distfield.py @@ -1,17 +1,21 @@ #!/usr/bin/env python3 -"""Model of the patched HashTable distance packing, to check the bit masks. +"""Model of HashTable's distance packing, at both compiled widths. -Mirrors HashTable::Convert / CalcDistAndType word-for-word so the masks can be -exercised without a C++ toolchain. If this fails, the C++ is wrong too. +Mirrors HashTable::Convert / CalcDistAndType mask-for-mask so the bit layout +can be exercised without a C++ toolchain. If this fails, the C++ is wrong too. + +DIST_WORDS = 2 is the default build (126-bit magnitude, 125-bit intervals); +DIST_WORDS = 4 is -DWIDE_DIST (254-bit magnitude, 253-bit intervals). The flags +live in the top two bits of the top word at either width, which is what lets +one implementation serve both. """ N_ORDER = 0xFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFEBAAEDCE6AF48A03BBFD25E8CD0364141 +NB64BLOCK = 5 -SIGN = 0x8000000000000000 # b255 of the 256-bit field, i.e. b63 of word 3 -TYPE = 0x4000000000000000 # b254 -MAG3 = 0x3FFFFFFFFFFFFFFF # b253..b192 - -MAX_INTERVAL_BITS = 253 +SIGN = 0x8000000000000000 # top bit of the top word +TYPE = 0x4000000000000000 +MAG = 0x3FFFFFFFFFFFFFFF class Overflow(Exception): @@ -22,94 +26,89 @@ def words(v, n): return [(v >> (64 * i)) & 0xFFFFFFFFFFFFFFFF for i in range(n)] -def convert256(d_signed, type_bit): - """HashTable::Convert, 256-bit path. d_signed is the true signed distance.""" - d = d_signed % N_ORDER # how Int holds it: mod n - w = words(d, 4) +def max_interval_bits(dw): + return 253 if dw == 4 else 125 + + +def convert(d_signed, type_bit, dw): + """HashTable::Convert -- pack a signed distance into DIST_WORDS words.""" + d = d_signed % N_ORDER + w = words(d, NB64BLOCK) sign = 0 if w[3] > 0x7FFFFFFFFFFFFFFF: # upper half means negative d = (-d_signed) % N_ORDER # ModNegK1order - w = words(d, 4) + w = words(d, NB64BLOCK) sign = SIGN - if w[3] & 0xC000000000000000: - raise Overflow("distance exceeds 254 bits") - return [w[0], w[1], w[2], (w[3] & MAG3) | sign | (type_bit << 62)] - - -def calc_dist_and_type256(D): - """HashTable::CalcDistAndType, 256-bit path.""" - ktype = 1 if (D[3] & TYPE) else 0 - sign = 1 if (D[3] & SIGN) else 0 - mag = D[0] | (D[1] << 64) | (D[2] << 128) | ((D[3] & MAG3) << 192) + # magnitude must fit in DIST_MAG_BITS: flag bits clear in the top word, + # and nothing at all above it + if w[dw - 1] & (SIGN | TYPE) or any(w[i] for i in range(dw, NB64BLOCK)): + raise Overflow("distance exceeds %d bits" % (64 * dw - 2)) + out = w[:dw] + out[dw - 1] = (out[dw - 1] & MAG) | sign | (type_bit << 62) + return out + + +def calc_dist_and_type(D, dw): + """HashTable::CalcDistAndType -- unpack it again.""" + ktype = 1 if (D[dw - 1] & TYPE) else 0 + sign = 1 if (D[dw - 1] & SIGN) else 0 + top = D[dw - 1] & MAG + mag = sum((top if i == dw - 1 else D[i]) << (64 * i) for i in range(dw)) return ((-mag) % N_ORDER if sign else mag), ktype -def convert128(d_signed, type_bit): - """Legacy 126-bit path, now guarded instead of truncating.""" - d = d_signed % N_ORDER - w = words(d, 4) - sign = 0 - if w[3] > 0x7FFFFFFFFFFFFFFF: - d = (-d_signed) % N_ORDER - w = words(d, 4) - sign = 1 << 63 - if w[3] or w[2] or (w[1] & 0xC000000000000000): - raise Overflow("distance exceeds 126 bits") - return [w[0], (w[1] & 0x3FFFFFFFFFFFFFFF) | sign | (type_bit << 62)] - - -def widen(D128): - """HashTable::Widen -- legacy 126-bit entry into the 254-bit field.""" - return [D128[0], D128[1] & 0x3FFFFFFFFFFFFFFF, 0, - D128[1] & 0xC000000000000000] - - -def roundtrip(d_signed, type_bit): - got, ktype = calc_dist_and_type256(convert256(d_signed, type_bit)) +def roundtrip(d_signed, type_bit, dw): + got, ktype = calc_dist_and_type(convert(d_signed, type_bit, dw), dw) return got == d_signed % N_ORDER and ktype == type_bit def demo(): - # Puzzle #140 lives in [2^139, 2^140); walks overshoot, so exercise well past it. - for bits in (139, 140, 160, 200, 253): - for t in (0, 1): - v = (1 << bits) - 12345 - assert roundtrip(v, t), (bits, t, "positive") - assert roundtrip(-v, t), (bits, t, "negative") - - # Boundary: 253 bits of magnitude fits, 254 does not. - assert roundtrip((1 << 253) - 1, 0) + for dw in (2, 4): + limit = max_interval_bits(dw) + + # Distances up to the documented interval limit survive, both signs, + # both kangaroo types. + for bits in (8, 64, 100, limit - 1, limit): + for t in (0, 1): + v = (1 << bits) - 12345 + assert roundtrip(v, t, dw), (dw, bits, t, "positive") + assert roundtrip(-v, t, dw), (dw, bits, t, "negative") + + # Exactly at the magnitude boundary is fine; one bit past is refused, + # never truncated. This is the whole point -- the original code masked + # here and returned a wrong key with no error. + mag_bits = 64 * dw - 2 + assert roundtrip((1 << mag_bits) - 1, 0, dw), dw + try: + convert(1 << mag_bits, 0, dw) + raise AssertionError("dw=%d: %d-bit magnitude must be rejected" + % (dw, mag_bits + 1)) + except Overflow: + pass + + # Flags must never collide with magnitude bits. + assert SIGN | TYPE | MAG == 0xFFFFFFFFFFFFFFFF + assert SIGN & MAG == 0 and TYPE & MAG == 0 and SIGN & TYPE == 0 + + # ENTRY layout: 16-byte x + 8*DIST_WORDS distance, no padding. + assert 16 + 8 * dw == (32 if dw == 2 else 48) + + # A puzzle #140 distance (2^139) is exactly what the default build must + # refuse and the wide build must accept. try: - convert256(1 << 254, 0) - raise AssertionError("254-bit magnitude must be rejected, not truncated") + convert(1 << 139, 0, 2) + raise AssertionError("default width must refuse a 139-bit distance") except Overflow: pass + assert roundtrip(1 << 139, 0, 4) + assert roundtrip(1 << 139, 1, 4) - # The original bug: a #140-scale distance silently lost its high bits in the - # 126-bit field. Now it raises instead of returning a wrong key. - try: - convert128(1 << 139, 0) - raise AssertionError("126-bit path must reject a 139-bit distance") - except Overflow: - pass - - # Legacy entries still decode correctly once widened. - for v in (1, 42, (1 << 125) - 1): - for t in (0, 1): - got, ktype = calc_dist_and_type256(widen(convert128(v, t))) - assert got == v and ktype == t, (v, t, got, ktype) - got, ktype = calc_dist_and_type256(widen(convert128(-v, t))) - assert got == (-v) % N_ORDER and ktype == t, (v, t) - - # Flags must never collide with magnitude bits. - assert SIGN | TYPE | MAG3 == 0xFFFFFFFFFFFFFFFF - assert SIGN & MAG3 == 0 and TYPE & MAG3 == 0 and SIGN & TYPE == 0 - - # ENTRY layout: 16-byte x + 32-byte d, no padding. - assert 16 + 32 == 48 + # Puzzle #120 (2^119) is inside both. + assert roundtrip(1 << 119, 0, 2) and roundtrip(1 << 119, 0, 4) - print("distance-field self-check OK (254-bit magnitude, max interval %d bits)" - % MAX_INTERVAL_BITS) + print("distance-field self-check OK") + print(" default magnitude 126 bits, intervals to 125") + print(" WIDE_DIST magnitude 254 bits, intervals to 253") if __name__ == "__main__":