Fix silent key corruption above 125-bit intervals: widen distance field to 256 bits

.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

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<int128_t>& kangs,Int* x,Int* y,Int* d) {
+void Kangaroo::FetchWalks(uint64_t nbWalk,std::vector<dist_t>& 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<int128_t> kangs;
+  vector<dist_t> kangs;
   if(saveKangarooByServer) {
     ::printf("FectchKangaroosFromServer");
     if(!GetKangaroosFromServer(workFile,kangs))
@@ -492,22 +492,22 @@ void Kangaroo::SaveWork(uint64_t totalCount,double totalTime,TH_PARAM *threads,i
     if(saveKangarooByServer) {
 
       ::printf("\nSaveWork (Kangaroo->Server): %s",fileName.c_str());
-      vector<int128_t> kangs;
+      vector<dist_t> 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);
           kangs.push_back(D);
         }
       }
       SendKangaroosToServer(fileName,kangs);
-      size = kangs.size()*16 + 16;
+      size = kangs.size()*(8*DIST_WORDS) + 16;
       goto end;
 
     } else {

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;

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

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]))

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<ITEM> &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<ITEM> &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);
 

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;