Reproducing the testcase

[solardiz]

Thanks for the testcase introduced in response to #4. However, the ethminer.exe -U -M 30000 command in test/result.log doesn't currently produce the debugging output also shown in there, and there doesn't appear to be any code in the repository to produce that output (or is there?) - can that please be added, perhaps as a debugging mode enabled via a command-line option (that would be included in the command given in test/result.log, so it'd be clear how to reproduce the result from just looking at that file)?

[ifdefelse]

I have a private branch with a bunch of extra printf's to generate that output. I'll get that cleaned up and checked in under #ifdef's.

We've also been asked to provides some unit test cases, which I'll do at the same time.

[solardiz]

Thanks. I've successfully used the now added test-vectors.md to validate https://github.com/solardiz/c-progpow at block 30k and 10M. I am getting the same values for "digest" that are specified in there. However, I don't know what it means by "result". Can this be clarified, please?

[solardiz]

Answering my own question from April 1, result is just the final Keccak hash of digest. I find it weird that README.md's progPowHash (pseudo-)code does compute this final Keccak hash yet drops its value and only returns digest. Perhaps instead of ending in:

    // keccak(header .. keccak(header..nonce) .. digest);
    keccak_f800_progpow(header, seed, digest);
    
    return digest;

it should end in:

    // keccak(header .. keccak(header..nonce) .. digest);
    return keccak_f800_progpow(header, seed, digest);

[solardiz]

I have a private branch with a bunch of extra printf's to generate that output. I'll get that cleaned up and checked in under #ifdef's.

Meanwhile, here's a hack of my own that I now use for a similar purpose:

diff --git a/libethash-cl/CLMiner.cpp b/libethash-cl/CLMiner.cpp
index 2a591f8..64a241f 100644
--- a/libethash-cl/CLMiner.cpp
+++ b/libethash-cl/CLMiner.cpp
@@ -320,7 +320,8 @@ void CLMiner::workLoop()
               assert(target > 0);
 
               // Update header constant buffer.
-              m_queue.enqueueWriteBuffer(m_header, CL_FALSE, 0, w.header.size, w.header.data());
+              //m_queue.enqueueWriteBuffer(m_header, CL_FALSE, 0, w.header.size, w.header.data());
+              m_queue.enqueueWriteBuffer(m_header, CL_FALSE, 0, 32, h256("ffeeddccbbaa9988776655443322110000112233445566778899aabbccddeeff").data());
               m_queue.enqueueWriteBuffer(m_searchBuffer, CL_FALSE, 0, sizeof(c_zero), &c_zero);
 
               m_searchKernel.setArg(0, m_searchBuffer);  // Supply output buffer to kernel.
@@ -335,6 +336,8 @@ void CLMiner::workLoop()
               else
                   startNonce = get_start_nonce();
 
+              startNonce = 0x123456789abcdef0ULL;
+
               clswitchlog << "Switch time"
                   << std::chrono::duration_cast<std::chrono::milliseconds>(std::chrono::high_resolution_clock::now() - workSwitchStart).count()
                   << "ms.";
diff --git a/libethash-cl/CLMiner_kernel.cl b/libethash-cl/CLMiner_kernel.cl
index 2214f96..6012278 100644
--- a/libethash-cl/CLMiner_kernel.cl
+++ b/libethash-cl/CLMiner_kernel.cl
@@ -221,6 +221,21 @@ __kernel void ethash_search(
             digest = digest_temp;
     }
 
+    if (gid == 0)
+    {
+        uint b[32];
+        for (int i = 0; i < 32; i++)
+            b[i] = (digest.uint32s[i / 4] >> (8 * (i % 4))) & 0xff;
+        printf("Nonce = %16lx Digest = "
+            "%02x%02x%02x%02x%02x%02x%02x%02x%02x%02x%02x%02x%02x%02x%02x%02x"
+            "%02x%02x%02x%02x%02x%02x%02x%02x%02x%02x%02x%02x%02x%02x%02x%02x\n",
+            nonce,
+            b[0], b[1], b[2], b[3], b[4], b[5], b[6], b[7],
+            b[8], b[9], b[10], b[11], b[12], b[13], b[14], b[15],
+            b[16], b[17], b[18], b[19], b[20], b[21], b[22], b[23],
+            b[24], b[25], b[26], b[27], b[28], b[29], b[30], b[31]);
+    }
+
     // keccak(header .. keccak(header..nonce) .. digest);
     if (keccak_f800(g_header, seed, digest) < target)
     {
diff --git a/libethash-cuda/CUDAMiner.cpp b/libethash-cuda/CUDAMiner.cpp
index 91c1655..5b1b431 100644
--- a/libethash-cuda/CUDAMiner.cpp
+++ b/libethash-cuda/CUDAMiner.cpp
@@ -555,6 +555,31 @@ void CUDAMiner::compileKernel(
   NVRTC_SAFE_CALL(nvrtcDestroyProgram(&prog));
 }
 
+static uint32_t bswap(uint32_t a)
+{
+  a = (a << 16) | (a >> 16);
+  return ((a & 0x00ff00ff) << 8) | ((a >> 8) & 0x00ff00ff);
+}
+
+static void unhex(hash32_t *dst, const char *src)
+{
+  const char *p = src;
+  uint32_t *q = dst->uint32s;
+  uint32_t v = 0;
+
+  while (*p && q <= &dst->uint32s[7]) {
+      if (*p >= '0' && *p <= '9')
+          v |= *p - '0';
+      else if (*p >= 'a' && *p <= 'f')
+          v |= *p - ('a' - 10);
+      else
+          break;
+      if (!((++p - src) & 7))
+          *q++ = bswap(v);
+      v <<= 4;
+  }
+}
+
 void CUDAMiner::search(
   uint8_t const* header,
   uint64_t target,
@@ -602,6 +627,13 @@ void CUDAMiner::search(
       }
   }
   const uint32_t batch_size = s_gridSize * s_blockSize;
+
+  m_starting_nonce = 0x123456789abcdef0ULL;
+  m_current_nonce = m_starting_nonce - batch_size;
+  unhex(&m_current_header, "ffeeddccbbaa9988776655443322110000112233445566778899aabbccddeeff");
+  m_current_target = -1;
+  s_noeval = true;
+
   while (true)
   {
       m_current_index++;
@@ -624,7 +656,11 @@ void CUDAMiner::search(
               for (unsigned int j = 0; j < found_count; j++) {
                   nonces[j] = nonce_base + buffer->result[j].gid;
                   if (s_noeval)
+                  {
                       memcpy(mixes[j].data(), (void *)&buffer->result[j].mix, sizeof(buffer->result[j].mix));
+                      if (nonces[j] == m_starting_nonce)
+                          cout << "Digest = " << mixes[j] << "\n";
+                  }
               }
           }
       }
diff --git a/libethash-cuda/CUDAMiner_cuda.h b/libethash-cuda/CUDAMiner_cuda.h
index a0629d5..573fa40 100644
--- a/libethash-cuda/CUDAMiner_cuda.h
+++ b/libethash-cuda/CUDAMiner_cuda.h
@@ -10,7 +10,7 @@
 // one solution per stream hash calculation
 // Leave room for up to 4 results. A power
 // of 2 here will yield better CUDA optimization
-#define SEARCH_RESULTS 4
+#define SEARCH_RESULTS 0x10000
 
 typedef struct {
   uint32_t count;
@@ -21,9 +21,10 @@ typedef struct {
   } result[SEARCH_RESULTS];
 } search_results;
 
-typedef struct
+typedef union
 {
   uint4 uint4s[32 / sizeof(uint4)];
+  uint32_t uint32s[32 / sizeof(uint32_t)];
 } hash32_t;
 
 typedef struct
diff --git a/libethash-cuda/CUDAMiner_kernel.cu b/libethash-cuda/CUDAMiner_kernel.cu
index 3f7666d..a33f14b 100644
--- a/libethash-cuda/CUDAMiner_kernel.cu
+++ b/libethash-cuda/CUDAMiner_kernel.cu
@@ -1,5 +1,5 @@
 #ifndef SEARCH_RESULTS
-#define SEARCH_RESULTS 4
+#define SEARCH_RESULTS 0x10000
 #endif
 
 typedef struct {