forked from ifdefelse/ProgPOW
-
Notifications
You must be signed in to change notification settings - Fork 0
/
progpow-debug.diff
147 lines (138 loc) · 4.64 KB
/
progpow-debug.diff
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
diff --git a/libethash-cl/CLMiner.cpp b/libethash-cl/CLMiner.cpp
index 9cc5961..9460342 100644
--- a/libethash-cl/CLMiner.cpp
+++ b/libethash-cl/CLMiner.cpp
@@ -322,7 +322,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.
@@ -337,6 +338,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 9a9b7e8..7c7f234 100644
--- a/libethash-cl/CLMiner_kernel.cl
+++ b/libethash-cl/CLMiner_kernel.cl
@@ -223,6 +223,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 59e837b..c4f8455 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..fe8a9a7 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 {