ProgPOW
ProgPOW copied to clipboard
Reproducing the testcase
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)?
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.
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?
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);
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 {