diff --git a/.gitignore b/.gitignore index b293acf6..765b6072 100644 --- a/.gitignore +++ b/.gitignore @@ -12,3 +12,20 @@ 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 + +# 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/Backup.cpp b/Backup.cpp index e90d74f4..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; - int128_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()*16 + 16; + size = kangs.size()*(8*DIST_WORDS) + 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/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 95a358a2..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][2]; + 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); - Add128(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 86be42b4..e27b64c4 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; @@ -364,11 +372,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]); } @@ -409,10 +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 + 10 * nbThreadPerGroup] = (uint64_t)NB_JUMP; + inputKangarooPinned[t + (8 + DIST_WORDS) * nbThreadPerGroup] = (uint64_t)NB_JUMP; #endif idx++; @@ -474,6 +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); @@ -528,6 +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 @@ -561,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 + 2*i,distance[i].bits64,16); - cudaMemcpyToSymbol(jD,jumpPinned,jumpSize/2); + 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)); @@ -654,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 + 12)); + it.kIdx = *((uint64_t*)(itemPtr + 8 + 2*DIST_WORDS)); uint64_t *x = (uint64_t *)itemPtr; it.x.bits64[0] = x[0]; @@ -666,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 8f61099e..6dcb3b7f 100644 --- a/GPU/GPUEngine.h +++ b/GPU/GPUEngine.h @@ -22,14 +22,18 @@ #include "../Constants.h" #include "../SECPK1/SECP256k1.h" +// 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 11 +#define KSIZE (9 + DIST_WORDS) #else -#define KSIZE 10 +#define KSIZE (8 + DIST_WORDS) #endif -#define ITEM_SIZE 56 -#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 b67e0055..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][2]; +__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]; @@ -122,6 +122,24 @@ __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]);} + +// 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) {\ USUBO(r[0],0ULL,r[0]); \ USUBC(r[1],0ULL,r[1]); \ @@ -170,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]; \ @@ -183,16 +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 *)idx)[0]; \ -out[pos*ITEM_SIZE32 + 14] = ((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][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][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][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][DIST_WORDS]) { #endif __syncthreads(); @@ -216,15 +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 + 10 * 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][2]) { +__device__ void LoadDists(uint64_t* a,uint64_t dist[GPU_GRP_SIZE][DIST_WORDS]) { __syncthreads(); @@ -235,6 +268,10 @@ __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]; +#if DIST_WORDS == 4 + d64[2] = (a)[IDX + 10 * blockDim.x + stride]; + d64[3] = (a)[IDX + 11 * blockDim.x + stride]; +#endif } @@ -271,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][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][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][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][DIST_WORDS]) { #endif __syncthreads(); @@ -296,9 +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 + 10 * blockDim.x + stride] = jumps[g]; + (a)[IDX + (8 + DIST_WORDS) * blockDim.x + stride] = jumps[g]; #endif } @@ -321,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][2]) { +__device__ void StoreDists(uint64_t* a,uint64_t dist[GPU_GRP_SIZE][DIST_WORDS]) { __syncthreads(); @@ -331,6 +372,10 @@ __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]; +#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 20ab658c..03d976fe 100644 --- a/HashTable.cpp +++ b/HashTable.cpp @@ -17,6 +17,7 @@ #include "HashTable.h" #include +#include #include #ifndef WIN64 #include @@ -54,17 +55,25 @@ uint64_t HashTable::GetNbItem() { } -ENTRY *HashTable::CreateEntry(int128_t *x,int128_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]; + for(int i = 0; i < DIST_WORDS; i++) + e->d.i64[i] = d->i64[i]; return e; } +bool HashTable::sameDist(dist_t *a,dist_t *b) { + + for(int i = 0; i < DIST_WORDS; i++) + if(a->i64[i] != b->i64[i]) return false; + return true; + +} + #define ADD_ENTRY(entry) { \ /* Shift the end of the index table */ \ for (int i = E[h].nbItem; i > st; i--) \ @@ -72,7 +81,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,dist_t *D) { uint64_t sign = 0; uint64_t type64 = (uint64_t)type << 62; @@ -80,28 +89,43 @@ 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; } - D->i64[1] |= sign; - D->i64[1] |= type64; + // 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); + } + + 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); } - -#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 +176,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 +189,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 +204,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 +234,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 +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; - int128_t D; + dist_t D; uint64_t h; Convert(x,d,type,&h,&X,&D); ENTRY* e = CreateEntry(&X,&D); @@ -239,22 +263,22 @@ 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,dist_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(dist_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[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]; + for(int i = 0; i < DIST_WORDS; i++) + kDist->bits64[i] = d.i64[i]; if(sign) kDist->ModNegK1order(); } @@ -287,7 +311,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 +405,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(dist_t),1,f); if(printPoint) { pointPrint++; if(pointPrint > point) { @@ -425,7 +449,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 +475,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(dist_t),1,f); E[h].items[i] = e; } diff --git a/HashTable.h b/HashTable.h index e59d898d..af30cb11 100644 --- a/HashTable.h +++ b/HashTable.h @@ -20,6 +20,7 @@ #include #include +#include "Constants.h" #include "SECPK1/Point.h" #ifdef WIN64 #include @@ -44,17 +45,50 @@ 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 -- 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 +#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) - int128_t d; // Travelled distance (b127=sign b126=kangaroo type, b125..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. 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 { uint32_t nbItem; @@ -69,7 +103,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,dist_t *d); int Add(uint64_t h,ENTRY *e); uint64_t GetNbItem(); void Reset(); @@ -88,15 +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,int128_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(int128_t d,Int* kDist,uint32_t* kType); + static void CalcDistAndType(dist_t d,Int* kDist,uint32_t* kType); private: - ENTRY *CreateEntry(int128_t *x,int128_t *d); + ENTRY *CreateEntry(int128_t *x,dist_t *d); static int compare(int128_t *i1,int128_t *i2); + static bool sameDist(dist_t *a,dist_t *b); std::string GetStr(int128_t *i); }; diff --git a/Kangaroo.cpp b/Kangaroo.cpp index 33104744..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,int128_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 70f5e342..c148563e 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; + dist_t d; } DP; @@ -117,9 +117,20 @@ typedef struct { } DP_CACHE; // Work file type +// 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) @@ -165,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,int128_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); @@ -181,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); @@ -204,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 e6fc10ff..b16b2027 100644 --- a/Network.cpp +++ b/Network.cpp @@ -44,7 +44,14 @@ static SOCKET serverSock = 0; #define SERVER_VERSION 3 +// 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 @@ -336,7 +343,7 @@ bool Kangaroo::HandleRequest(TH_PARAM *p) { uint64_t nbKangaroo = 0; uint32_t strSize; char fileName[256]; - int128_t* KBuff; + dist_t* KBuff; uint32_t nbK; uint32_t header = HEADKS; uint32_t version = 0; @@ -378,7 +385,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 = (dist_t*)malloc(KANG_PER_BLOCK * sizeof(dist_t)); while(nbKangaroo > 0) { @@ -389,15 +396,15 @@ bool Kangaroo::HandleRequest(TH_PARAM *p) { } for(uint32_t k = 0; k < nbK; k++) { - ::fread(&KBuff[k],16,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]; + for(int w = 0; w < DIST_WORDS; w++) + K.bits64[w] = KBuff[k].i64[w]; checkSum.Add(&K); } - PUTFREE("packet",p->clientSock,KBuff,nbK * 16,ntimeout,KBuff); + PUTFREE("packet",p->clientSock,KBuff,nbK * (8*DIST_WORDS),ntimeout,KBuff); nbKangaroo -= nbK; @@ -405,6 +412,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); @@ -422,7 +432,7 @@ bool Kangaroo::HandleRequest(TH_PARAM *p) { uint32_t fileNameSize; char fileNameTmp[264]; char fileName[256]; - int128_t *KBuff; + dist_t *KBuff; uint32_t nbK; uint32_t header = HEADKS; uint32_t version = 0; @@ -459,7 +469,7 @@ bool Kangaroo::HandleRequest(TH_PARAM *p) { checkSum.SetInt32(0); - KBuff = (int128_t *)malloc(KANG_PER_BLOCK*sizeof(int128_t)); + KBuff = (dist_t *)malloc(KANG_PER_BLOCK*sizeof(dist_t)); while(nbKangaroo>0) { @@ -469,14 +479,14 @@ 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 * (8*DIST_WORDS),ntimeout,KBuff); for(uint32_t k = 0; k < nbK; k++) { - ::fwrite(&KBuff[k],16,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]; + for(int w = 0; w < DIST_WORDS; w++) + K.bits64[w] = KBuff[k].i64[w]; checkSum.Add(&K); } @@ -490,6 +500,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 { @@ -696,7 +709,7 @@ void Kangaroo::RunServer() { } SetDP(initDPSize); - if(sizeof(DP) != 40) { + if(sizeof(DP) != (int)(8 + 16 + (8*DIST_WORDS))) { ::printf("Error: Invalid DP size struct\n"); exit(-1); } @@ -980,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; - int128_t* KBuff; + dist_t* KBuff; Int checkSum; WaitForServer(); @@ -1011,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; - int128_t *KBuff; + dist_t *KBuff; Int checkSum; WaitForServer(); @@ -1088,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 = (int128_t*)malloc(KANG_PER_BLOCK * sizeof(int128_t)); + KBuff = (dist_t*)malloc(KANG_PER_BLOCK * sizeof(dist_t)); checkSum.SetInt32(0); pos = 0; @@ -1107,17 +1123,17 @@ 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],sizeof(dist_t)); pos++; // Checksum Int K; K.SetInt32(0); - K.bits64[1] = KBuff[k].i64[1]; - K.bits64[0] = KBuff[k].i64[0]; + for(int w = 0; w < DIST_WORDS; w++) + K.bits64[w] = KBuff[k].i64[w]; checkSum.Add(&K); } - PUTFREE("packet",serverConn,KBuff,nbK * 16,ntimeout,KBuff); + PUTFREE("packet",serverConn,KBuff,nbK * (8*DIST_WORDS),ntimeout,KBuff); nbKangaroo -= nbK; @@ -1125,6 +1141,9 @@ bool Kangaroo::SendKangaroosToServer(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); } @@ -1154,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]; + 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_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..6a6cff67 --- /dev/null +++ b/build_gpu.bat @@ -0,0 +1,63 @@ +@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 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" + +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) + +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 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 %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 %WIDE% -D_CRT_SECURE_NO_WARNINGS ^ + -gencode=arch=compute_%CCAP%,code=sm_%CCAP% ^ + -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 %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 /Fo%OBJD%\ +if errorlevel 1 (echo COMPILE FAILED & exit /b 1) + +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/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/test140.txt b/test140.txt new file mode 100644 index 00000000..d47d9757 --- /dev/null +++ b/test140.txt @@ -0,0 +1,3 @@ +80000000000000000000000000000000000 +FFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFF +031F6A332D3C5C4F2DE2378C012F429CD109BA07D69690C6C701B6BB87860D6640 diff --git a/tools/check_layout.py b/tools/check_layout.py new file mode 100644 index 00000000..7bc8e3b6 --- /dev/null +++ b/tools/check_layout.py @@ -0,0 +1,228 @@ +#!/usr/bin/env python3 +"""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 +from pathlib import Path + +_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 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 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") + 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}")].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(dw): + """Device kangaroo record: px[4] py[4] dist[dw] (+lastJump) must fit KSIZE.""" + eng_h = read("GPU/GPUEngine.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_add_carry(dw): + """AddDist must select a carry chain covering exactly DIST_WORDS words.""" + math_h = read("GPU/GPUMath.h") + 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) + + # 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") + 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" + 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_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") + 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") + + 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__": + demo() 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/test_distfield.py b/tools/test_distfield.py new file mode 100644 index 00000000..bc33d09a --- /dev/null +++ b/tools/test_distfield.py @@ -0,0 +1,115 @@ +#!/usr/bin/env python3 +"""Model of HashTable's distance packing, at both compiled widths. + +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 # top bit of the top word +TYPE = 0x4000000000000000 +MAG = 0x3FFFFFFFFFFFFFFF + + +class Overflow(Exception): + pass + + +def words(v, n): + return [(v >> (64 * i)) & 0xFFFFFFFFFFFFFFFF for i in range(n)] + + +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, NB64BLOCK) + sign = SIGN + # 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 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(): + 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: + 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) + + # 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") + print(" default magnitude 126 bits, intervals to 125") + print(" WIDE_DIST magnitude 254 bits, intervals to 253") + + +if __name__ == "__main__": + demo() 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)