Skip to content

Commit 0659d69

Browse files
Christian BuchnerChristian Buchner
Christian Buchner
authored and
Christian Buchner
committed
Revision 0.6 with myriad-groestl and jackpot coin
1 parent 9efd648 commit 0659d69

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

55 files changed

+21454
-21350
lines changed

‎JHA/cuda_jha_keccak512.cu

+572-572
Large diffs are not rendered by default.

‎JHA/jackpotcoin.cu

+173-173
Original file line numberDiff line numberDiff line change
@@ -1,173 +1,173 @@
1-
2-
extern "C"
3-
{
4-
#include "sph/sph_keccak.h"
5-
#include "sph/sph_blake.h"
6-
#include "sph/sph_groestl.h"
7-
#include "sph/sph_jh.h"
8-
#include "sph/sph_skein.h"
9-
}
10-
11-
#include "miner.h"
12-
#include <stdint.h>
13-
14-
// aus cpu-miner.c
15-
extern int device_map[8];
16-
extern bool opt_benchmark;
17-
18-
// Speicher für Input/Output der verketteten Hashfunktionen
19-
static uint32_t *d_hash[8];
20-
21-
extern void jackpot_keccak512_cpu_init(int thr_id, int threads);
22-
extern void jackpot_keccak512_cpu_setBlock_88(void *pdata);
23-
extern void jackpot_keccak512_cpu_hash_88(int thr_id, int threads, uint32_t startNounce, uint32_t *d_hash, int order);
24-
25-
extern void quark_check_cpu_init(int thr_id, int threads);
26-
extern void quark_check_cpu_setTarget(const void *ptarget);
27-
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);
28-
29-
// Original jackpothash Funktion aus einem miner Quelltext
30-
inline unsigned int jackpothash(void *state, const void *input)
31-
{
32-
sph_blake512_context ctx_blake;
33-
sph_groestl512_context ctx_groestl;
34-
sph_jh512_context ctx_jh;
35-
sph_keccak512_context ctx_keccak;
36-
sph_skein512_context ctx_skein;
37-
38-
uint32_t hash[16];
39-
40-
sph_keccak512_init(&ctx_keccak);
41-
sph_keccak512 (&ctx_keccak, input, 88);
42-
sph_keccak512_close(&ctx_keccak, hash);
43-
44-
unsigned int round_mask = (
45-
(unsigned int)(((unsigned char *)input)[84]) << 0 |
46-
(unsigned int)(((unsigned char *)input)[85]) << 8 |
47-
(unsigned int)(((unsigned char *)input)[86]) << 16 |
48-
(unsigned int)(((unsigned char *)input)[87]) << 24 );
49-
unsigned int round_max = hash[0] & round_mask;
50-
unsigned int round;
51-
for (round = 0; round < round_max; round++) {
52-
switch (hash[0] & 3) {
53-
case 0:
54-
sph_blake512_init(&ctx_blake);
55-
sph_blake512 (&ctx_blake, hash, 64);
56-
sph_blake512_close(&ctx_blake, hash);
57-
break;
58-
case 1:
59-
sph_groestl512_init(&ctx_groestl);
60-
sph_groestl512 (&ctx_groestl, hash, 64);
61-
sph_groestl512_close(&ctx_groestl, hash);
62-
break;
63-
case 2:
64-
sph_jh512_init(&ctx_jh);
65-
sph_jh512 (&ctx_jh, hash, 64);
66-
sph_jh512_close(&ctx_jh, hash);
67-
break;
68-
case 3:
69-
sph_skein512_init(&ctx_skein);
70-
sph_skein512 (&ctx_skein, hash, 64);
71-
sph_skein512_close(&ctx_skein, hash);
72-
break;
73-
}
74-
}
75-
memcpy(state, hash, 32);
76-
77-
return round_max;
78-
}
79-
80-
81-
static int bit_population(uint32_t n){
82-
int c =0;
83-
while(n){
84-
c += n&1;
85-
n = n>>1;
86-
}
87-
return c;
88-
}
89-
90-
extern "C" int scanhash_jackpot(int thr_id, uint32_t *pdata,
91-
const uint32_t *ptarget, uint32_t max_nonce,
92-
unsigned long *hashes_done)
93-
{
94-
const uint32_t first_nonce = pdata[19];
95-
96-
// TODO: entfernen für eine Release! Ist nur zum Testen!
97-
if (opt_benchmark) {
98-
((uint32_t*)ptarget)[7] = 0x00000f;
99-
((uint32_t*)pdata)[21] = 0x07000000; // round_mask von 7 vorgeben
100-
}
101-
102-
const uint32_t Htarg = ptarget[7];
103-
104-
const int throughput = 256*4096; // 100;
105-
106-
static bool init[8] = {0,0,0,0,0,0,0,0};
107-
if (!init[thr_id])
108-
{
109-
cudaSetDevice(device_map[thr_id]);
110-
111-
// Konstanten kopieren, Speicher belegen
112-
cudaMalloc(&d_hash[thr_id], 16 * sizeof(uint32_t) * throughput);
113-
jackpot_keccak512_cpu_init(thr_id, throughput);
114-
quark_check_cpu_init(thr_id, throughput);
115-
init[thr_id] = true;
116-
}
117-
118-
uint32_t endiandata[22];
119-
for (int k=0; k < 22; k++)
120-
be32enc(&endiandata[k], ((uint32_t*)pdata)[k]);
121-
122-
unsigned int round_mask = (
123-
(unsigned int)(((unsigned char *)endiandata)[84]) << 0 |
124-
(unsigned int)(((unsigned char *)endiandata)[85]) << 8 |
125-
(unsigned int)(((unsigned char *)endiandata)[86]) << 16 |
126-
(unsigned int)(((unsigned char *)endiandata)[87]) << 24 );
127-
128-
// Zählen wie viele Bits in round_mask gesetzt sind
129-
int bitcount = bit_population(round_mask);
130-
131-
jackpot_keccak512_cpu_setBlock_88((void*)endiandata);
132-
quark_check_cpu_setTarget(ptarget);
133-
134-
do {
135-
int order = 0;
136-
137-
// erstes Blake512 Hash mit CUDA
138-
jackpot_keccak512_cpu_hash_88(thr_id, throughput, pdata[19], d_hash[thr_id], order++);
139-
140-
// TODO: hier fehlen jetzt natürlich noch die anderen Hashrunden.
141-
// bei round_mask=7 haben wir eine 1:8 Chance, dass das Hash dennoch
142-
// die Kriterien erfüllt wenn hash[0] & round_mask zufällig 0 ist.
143-
144-
// Scan nach Gewinner Hashes auf der GPU
145-
uint32_t foundNonce = quark_check_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
146-
if (foundNonce != 0xffffffff)
147-
{
148-
uint32_t vhash64[8];
149-
be32enc(&endiandata[19], foundNonce);
150-
151-
// diese jackpothash Funktion gibt die Zahl der zusätzlichen Runden zurück
152-
unsigned int rounds = jackpothash(vhash64, endiandata);
153-
154-
// wir akzeptieren nur solche Hashes wo ausschliesslich Keccak verwendet wurde
155-
if (rounds == 0) {
156-
if ((vhash64[7]<=Htarg) && fulltest(vhash64, ptarget)) {
157-
158-
pdata[19] = foundNonce;
159-
*hashes_done = (foundNonce - first_nonce + 1) / (1 << bitcount);
160-
return 1;
161-
} else {
162-
applog(LOG_INFO, "GPU #%d: result for nonce $%08X does not validate on CPU (%d rounds)!", thr_id, foundNonce, rounds);
163-
}
164-
}
165-
}
166-
167-
pdata[19] += throughput;
168-
169-
} while (pdata[19] < max_nonce && !work_restart[thr_id].restart);
170-
171-
*hashes_done = (pdata[19] - first_nonce + 1) / (1 << bitcount);
172-
return 0;
173-
}
1+
2+
extern "C"
3+
{
4+
#include "sph/sph_keccak.h"
5+
#include "sph/sph_blake.h"
6+
#include "sph/sph_groestl.h"
7+
#include "sph/sph_jh.h"
8+
#include "sph/sph_skein.h"
9+
}
10+
11+
#include "miner.h"
12+
#include <stdint.h>
13+
14+
// aus cpu-miner.c
15+
extern int device_map[8];
16+
extern bool opt_benchmark;
17+
18+
// Speicher für Input/Output der verketteten Hashfunktionen
19+
static uint32_t *d_hash[8];
20+
21+
extern void jackpot_keccak512_cpu_init(int thr_id, int threads);
22+
extern void jackpot_keccak512_cpu_setBlock_88(void *pdata);
23+
extern void jackpot_keccak512_cpu_hash_88(int thr_id, int threads, uint32_t startNounce, uint32_t *d_hash, int order);
24+
25+
extern void quark_check_cpu_init(int thr_id, int threads);
26+
extern void quark_check_cpu_setTarget(const void *ptarget);
27+
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);
28+
29+
// Original jackpothash Funktion aus einem miner Quelltext
30+
inline unsigned int jackpothash(void *state, const void *input)
31+
{
32+
sph_blake512_context ctx_blake;
33+
sph_groestl512_context ctx_groestl;
34+
sph_jh512_context ctx_jh;
35+
sph_keccak512_context ctx_keccak;
36+
sph_skein512_context ctx_skein;
37+
38+
uint32_t hash[16];
39+
40+
sph_keccak512_init(&ctx_keccak);
41+
sph_keccak512 (&ctx_keccak, input, 88);
42+
sph_keccak512_close(&ctx_keccak, hash);
43+
44+
unsigned int round_mask = (
45+
(unsigned int)(((unsigned char *)input)[84]) << 0 |
46+
(unsigned int)(((unsigned char *)input)[85]) << 8 |
47+
(unsigned int)(((unsigned char *)input)[86]) << 16 |
48+
(unsigned int)(((unsigned char *)input)[87]) << 24 );
49+
unsigned int round_max = hash[0] & round_mask;
50+
unsigned int round;
51+
for (round = 0; round < round_max; round++) {
52+
switch (hash[0] & 3) {
53+
case 0:
54+
sph_blake512_init(&ctx_blake);
55+
sph_blake512 (&ctx_blake, hash, 64);
56+
sph_blake512_close(&ctx_blake, hash);
57+
break;
58+
case 1:
59+
sph_groestl512_init(&ctx_groestl);
60+
sph_groestl512 (&ctx_groestl, hash, 64);
61+
sph_groestl512_close(&ctx_groestl, hash);
62+
break;
63+
case 2:
64+
sph_jh512_init(&ctx_jh);
65+
sph_jh512 (&ctx_jh, hash, 64);
66+
sph_jh512_close(&ctx_jh, hash);
67+
break;
68+
case 3:
69+
sph_skein512_init(&ctx_skein);
70+
sph_skein512 (&ctx_skein, hash, 64);
71+
sph_skein512_close(&ctx_skein, hash);
72+
break;
73+
}
74+
}
75+
memcpy(state, hash, 32);
76+
77+
return round_max;
78+
}
79+
80+
81+
static int bit_population(uint32_t n){
82+
int c =0;
83+
while(n){
84+
c += n&1;
85+
n = n>>1;
86+
}
87+
return c;
88+
}
89+
90+
extern "C" int scanhash_jackpot(int thr_id, uint32_t *pdata,
91+
const uint32_t *ptarget, uint32_t max_nonce,
92+
unsigned long *hashes_done)
93+
{
94+
const uint32_t first_nonce = pdata[19];
95+
96+
// TODO: entfernen für eine Release! Ist nur zum Testen!
97+
if (opt_benchmark) {
98+
((uint32_t*)ptarget)[7] = 0x00000f;
99+
((uint32_t*)pdata)[21] = 0x07000000; // round_mask von 7 vorgeben
100+
}
101+
102+
const uint32_t Htarg = ptarget[7];
103+
104+
const int throughput = 256*4096; // 100;
105+
106+
static bool init[8] = {0,0,0,0,0,0,0,0};
107+
if (!init[thr_id])
108+
{
109+
cudaSetDevice(device_map[thr_id]);
110+
111+
// Konstanten kopieren, Speicher belegen
112+
cudaMalloc(&d_hash[thr_id], 16 * sizeof(uint32_t) * throughput);
113+
jackpot_keccak512_cpu_init(thr_id, throughput);
114+
quark_check_cpu_init(thr_id, throughput);
115+
init[thr_id] = true;
116+
}
117+
118+
uint32_t endiandata[22];
119+
for (int k=0; k < 22; k++)
120+
be32enc(&endiandata[k], ((uint32_t*)pdata)[k]);
121+
122+
unsigned int round_mask = (
123+
(unsigned int)(((unsigned char *)endiandata)[84]) << 0 |
124+
(unsigned int)(((unsigned char *)endiandata)[85]) << 8 |
125+
(unsigned int)(((unsigned char *)endiandata)[86]) << 16 |
126+
(unsigned int)(((unsigned char *)endiandata)[87]) << 24 );
127+
128+
// Zählen wie viele Bits in round_mask gesetzt sind
129+
int bitcount = bit_population(round_mask);
130+
131+
jackpot_keccak512_cpu_setBlock_88((void*)endiandata);
132+
quark_check_cpu_setTarget(ptarget);
133+
134+
do {
135+
int order = 0;
136+
137+
// erstes Blake512 Hash mit CUDA
138+
jackpot_keccak512_cpu_hash_88(thr_id, throughput, pdata[19], d_hash[thr_id], order++);
139+
140+
// TODO: hier fehlen jetzt natürlich noch die anderen Hashrunden.
141+
// bei round_mask=7 haben wir eine 1:8 Chance, dass das Hash dennoch
142+
// die Kriterien erfüllt wenn hash[0] & round_mask zufällig 0 ist.
143+
144+
// Scan nach Gewinner Hashes auf der GPU
145+
uint32_t foundNonce = quark_check_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
146+
if (foundNonce != 0xffffffff)
147+
{
148+
uint32_t vhash64[8];
149+
be32enc(&endiandata[19], foundNonce);
150+
151+
// diese jackpothash Funktion gibt die Zahl der zusätzlichen Runden zurück
152+
unsigned int rounds = jackpothash(vhash64, endiandata);
153+
154+
// wir akzeptieren nur solche Hashes wo ausschliesslich Keccak verwendet wurde
155+
if (rounds == 0) {
156+
if ((vhash64[7]<=Htarg) && fulltest(vhash64, ptarget)) {
157+
158+
pdata[19] = foundNonce;
159+
*hashes_done = (foundNonce - first_nonce + 1) / (1 << bitcount);
160+
return 1;
161+
} else {
162+
applog(LOG_INFO, "GPU #%d: result for nonce $%08X does not validate on CPU (%d rounds)!", thr_id, foundNonce, rounds);
163+
}
164+
}
165+
}
166+
167+
pdata[19] += throughput;
168+
169+
} while (pdata[19] < max_nonce && !work_restart[thr_id].restart);
170+
171+
*hashes_done = (pdata[19] - first_nonce + 1) / (1 << bitcount);
172+
return 0;
173+
}

0 commit comments

Comments
 (0)