-
Notifications
You must be signed in to change notification settings - Fork 1.7k
/
jackpotcoin.cu
173 lines (140 loc) · 5.31 KB
/
jackpotcoin.cu
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
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
extern "C"
{
#include "sph/sph_keccak.h"
#include "sph/sph_blake.h"
#include "sph/sph_groestl.h"
#include "sph/sph_jh.h"
#include "sph/sph_skein.h"
}
#include "miner.h"
#include <stdint.h>
// aus cpu-miner.c
extern int device_map[8];
extern bool opt_benchmark;
// Speicher für Input/Output der verketteten Hashfunktionen
static uint32_t *d_hash[8];
extern void jackpot_keccak512_cpu_init(int thr_id, int threads);
extern void jackpot_keccak512_cpu_setBlock_88(void *pdata);
extern void jackpot_keccak512_cpu_hash_88(int thr_id, int threads, uint32_t startNounce, uint32_t *d_hash, int order);
extern void quark_check_cpu_init(int thr_id, int threads);
extern void quark_check_cpu_setTarget(const void *ptarget);
extern uint32_t quark_check_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_inputHash, int order);
// Original jackpothash Funktion aus einem miner Quelltext
inline unsigned int jackpothash(void *state, const void *input)
{
sph_blake512_context ctx_blake;
sph_groestl512_context ctx_groestl;
sph_jh512_context ctx_jh;
sph_keccak512_context ctx_keccak;
sph_skein512_context ctx_skein;
uint32_t hash[16];
sph_keccak512_init(&ctx_keccak);
sph_keccak512 (&ctx_keccak, input, 88);
sph_keccak512_close(&ctx_keccak, hash);
unsigned int round_mask = (
(unsigned int)(((unsigned char *)input)[84]) << 0 |
(unsigned int)(((unsigned char *)input)[85]) << 8 |
(unsigned int)(((unsigned char *)input)[86]) << 16 |
(unsigned int)(((unsigned char *)input)[87]) << 24 );
unsigned int round_max = hash[0] & round_mask;
unsigned int round;
for (round = 0; round < round_max; round++) {
switch (hash[0] & 3) {
case 0:
sph_blake512_init(&ctx_blake);
sph_blake512 (&ctx_blake, hash, 64);
sph_blake512_close(&ctx_blake, hash);
break;
case 1:
sph_groestl512_init(&ctx_groestl);
sph_groestl512 (&ctx_groestl, hash, 64);
sph_groestl512_close(&ctx_groestl, hash);
break;
case 2:
sph_jh512_init(&ctx_jh);
sph_jh512 (&ctx_jh, hash, 64);
sph_jh512_close(&ctx_jh, hash);
break;
case 3:
sph_skein512_init(&ctx_skein);
sph_skein512 (&ctx_skein, hash, 64);
sph_skein512_close(&ctx_skein, hash);
break;
}
}
memcpy(state, hash, 32);
return round_max;
}
static int bit_population(uint32_t n){
int c =0;
while(n){
c += n&1;
n = n>>1;
}
return c;
}
extern "C" int scanhash_jackpot(int thr_id, uint32_t *pdata,
const uint32_t *ptarget, uint32_t max_nonce,
unsigned long *hashes_done)
{
const uint32_t first_nonce = pdata[19];
// TODO: entfernen für eine Release! Ist nur zum Testen!
if (opt_benchmark) {
((uint32_t*)ptarget)[7] = 0x00000f;
((uint32_t*)pdata)[21] = 0x07000000; // round_mask von 7 vorgeben
}
const uint32_t Htarg = ptarget[7];
const int throughput = 256*4096; // 100;
static bool init[8] = {0,0,0,0,0,0,0,0};
if (!init[thr_id])
{
cudaSetDevice(device_map[thr_id]);
// Konstanten kopieren, Speicher belegen
cudaMalloc(&d_hash[thr_id], 16 * sizeof(uint32_t) * throughput);
jackpot_keccak512_cpu_init(thr_id, throughput);
quark_check_cpu_init(thr_id, throughput);
init[thr_id] = true;
}
uint32_t endiandata[22];
for (int k=0; k < 22; k++)
be32enc(&endiandata[k], ((uint32_t*)pdata)[k]);
unsigned int round_mask = (
(unsigned int)(((unsigned char *)endiandata)[84]) << 0 |
(unsigned int)(((unsigned char *)endiandata)[85]) << 8 |
(unsigned int)(((unsigned char *)endiandata)[86]) << 16 |
(unsigned int)(((unsigned char *)endiandata)[87]) << 24 );
// Zählen wie viele Bits in round_mask gesetzt sind
int bitcount = bit_population(round_mask);
jackpot_keccak512_cpu_setBlock_88((void*)endiandata);
quark_check_cpu_setTarget(ptarget);
do {
int order = 0;
// erstes Blake512 Hash mit CUDA
jackpot_keccak512_cpu_hash_88(thr_id, throughput, pdata[19], d_hash[thr_id], order++);
// TODO: hier fehlen jetzt natürlich noch die anderen Hashrunden.
// bei round_mask=7 haben wir eine 1:8 Chance, dass das Hash dennoch
// die Kriterien erfüllt wenn hash[0] & round_mask zufällig 0 ist.
// Scan nach Gewinner Hashes auf der GPU
uint32_t foundNonce = quark_check_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
if (foundNonce != 0xffffffff)
{
uint32_t vhash64[8];
be32enc(&endiandata[19], foundNonce);
// diese jackpothash Funktion gibt die Zahl der zusätzlichen Runden zurück
unsigned int rounds = jackpothash(vhash64, endiandata);
// wir akzeptieren nur solche Hashes wo ausschliesslich Keccak verwendet wurde
if (rounds == 0) {
if ((vhash64[7]<=Htarg) && fulltest(vhash64, ptarget)) {
pdata[19] = foundNonce;
*hashes_done = (foundNonce - first_nonce + 1) / (1 << bitcount);
return 1;
} else {
applog(LOG_INFO, "GPU #%d: result for nonce $%08X does not validate on CPU (%d rounds)!", thr_id, foundNonce, rounds);
}
}
}
pdata[19] += throughput;
} while (pdata[19] < max_nonce && !work_restart[thr_id].restart);
*hashes_done = (pdata[19] - first_nonce + 1) / (1 << bitcount);
return 0;
}