Skip to content

Commit ea177b7

Browse files
JeanLucPonsJeanLucPons
committedMay 13, 2019
Added wildcard search
·
1.191.14
1 parent 8f53925 commit ea177b7

17 files changed

+1428
-349
lines changed
 

‎GPU/GPUBase58.h

+130
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,130 @@
1+
/*
2+
* This file is part of the VanitySearch distribution (https://github.com/JeanLucPons/VanitySearch).
3+
* Copyright (c) 2019 Jean Luc PONS.
4+
*
5+
* This program is free software: you can redistribute it and/or modify
6+
* it under the terms of the GNU General Public License as published by
7+
* the Free Software Foundation, version 3.
8+
*
9+
* This program is distributed in the hope that it will be useful, but
10+
* WITHOUT ANY WARRANTY; without even the implied warranty of
11+
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
12+
* General Public License for more details.
13+
*
14+
* You should have received a copy of the GNU General Public License
15+
* along with this program. If not, see <http://www.gnu.org/licenses/>.
16+
*/
17+
18+
// ---------------------------------------------------------------------------------
19+
// Base58
20+
// ---------------------------------------------------------------------------------
21+
22+
__device__ __constant__ char *pszBase58 = (char *)"123456789ABCDEFGHJKLMNPQRSTUVWXYZabcdefghijkmnopqrstuvwxyz";
23+
24+
__device__ __constant__ int8_t b58digits_map[] = {
25+
-1,-1,-1,-1,-1,-1,-1,-1, -1,-1,-1,-1,-1,-1,-1,-1,
26+
-1,-1,-1,-1,-1,-1,-1,-1, -1,-1,-1,-1,-1,-1,-1,-1,
27+
-1,-1,-1,-1,-1,-1,-1,-1, -1,-1,-1,-1,-1,-1,-1,-1,
28+
-1, 0, 1, 2, 3, 4, 5, 6, 7, 8,-1,-1,-1,-1,-1,-1,
29+
-1, 9,10,11,12,13,14,15, 16,-1,17,18,19,20,21,-1,
30+
22,23,24,25,26,27,28,29, 30,31,32,-1,-1,-1,-1,-1,
31+
-1,33,34,35,36,37,38,39, 40,41,42,43,-1,44,45,46,
32+
47,48,49,50,51,52,53,54, 55,56,57,-1,-1,-1,-1,-1,
33+
};
34+
35+
__device__ __noinline__ void _GetAddress(int type,uint32_t *hash,char *b58Add) {
36+
37+
uint32_t addBytes[16];
38+
uint32_t s[16];
39+
unsigned char A[25];
40+
unsigned char *addPtr = A;
41+
int retPos = 0;
42+
unsigned char digits[128];
43+
44+
switch (type) {
45+
46+
case P2PKH:
47+
A[0] = 0x00;
48+
break;
49+
50+
case P2SH:
51+
A[0] = 0x05;
52+
break;
53+
54+
}
55+
memcpy(A + 1, (char *)hash, 20);
56+
57+
// Compute checksum
58+
59+
addBytes[0] = __byte_perm(hash[0], (uint32_t)A[0], 0x4012);
60+
addBytes[1] = __byte_perm(hash[0], hash[1], 0x3456);
61+
addBytes[2] = __byte_perm(hash[1], hash[2], 0x3456);
62+
addBytes[3] = __byte_perm(hash[2], hash[3], 0x3456);
63+
addBytes[4] = __byte_perm(hash[3], hash[4], 0x3456);
64+
addBytes[5] = __byte_perm(hash[4], 0x80, 0x3456);
65+
addBytes[6] = 0;
66+
addBytes[7] = 0;
67+
addBytes[8] = 0;
68+
addBytes[9] = 0;
69+
addBytes[10] = 0;
70+
addBytes[11] = 0;
71+
addBytes[12] = 0;
72+
addBytes[13] = 0;
73+
addBytes[14] = 0;
74+
addBytes[15] = 0xA8;
75+
76+
SHA256Initialize(s);
77+
SHA256Transform(s, addBytes);
78+
79+
#pragma unroll 8
80+
for (int i = 0; i < 8; i++)
81+
addBytes[i] = s[i];
82+
83+
addBytes[8] = 0x80000000;
84+
addBytes[9] = 0;
85+
addBytes[10] = 0;
86+
addBytes[11] = 0;
87+
addBytes[12] = 0;
88+
addBytes[13] = 0;
89+
addBytes[14] = 0;
90+
addBytes[15] = 0x100;
91+
92+
SHA256Initialize(s);
93+
SHA256Transform(s, addBytes);
94+
95+
A[21] = ((uint8_t *)s)[3];
96+
A[22] = ((uint8_t *)s)[2];
97+
A[23] = ((uint8_t *)s)[1];
98+
A[24] = ((uint8_t *)s)[0];
99+
100+
// Base58
101+
102+
// Skip leading zeroes
103+
while(addPtr[0] == 0) {
104+
b58Add[retPos++] = '1';
105+
addPtr++;
106+
}
107+
int length = 25-retPos;
108+
109+
int digitslen = 1;
110+
digits[0] = 0;
111+
for (int i = 0; i < length; i++) {
112+
uint32_t carry = addPtr[i];
113+
for (int j = 0; j < digitslen; j++) {
114+
carry += (uint32_t)(digits[j]) << 8;
115+
digits[j] = (unsigned char)(carry % 58);
116+
carry /= 58;
117+
}
118+
while (carry > 0) {
119+
digits[digitslen++] = (unsigned char)(carry % 58);
120+
carry /= 58;
121+
}
122+
}
123+
124+
// reverse
125+
for (int i = 0; i < digitslen; i++)
126+
b58Add[retPos++] = (pszBase58[digits[digitslen - 1 - i]]);
127+
128+
b58Add[retPos] = 0;
129+
130+
}

‎GPU/GPUCompute.h

+79-49
Original file line numberDiff line numberDiff line change
@@ -24,7 +24,7 @@
2424
// We use affine coordinates for elliptic curve point (ie Z=1)
2525

2626
__device__ __noinline__ void CheckPoint(uint32_t *_h, int32_t incr, int32_t endo, int32_t mode,prefix_t *prefix,
27-
uint32_t *lookup32, uint32_t maxFound, uint32_t *out) {
27+
uint32_t *lookup32, uint32_t maxFound, uint32_t *out,int type) {
2828

2929
uint32_t off;
3030
prefixl_t l32;
@@ -36,43 +36,59 @@ __device__ __noinline__ void CheckPoint(uint32_t *_h, int32_t incr, int32_t endo
3636
uint32_t mi;
3737
uint32_t lmi;
3838
uint32_t tid = (blockIdx.x*blockDim.x) + threadIdx.x;
39-
40-
pr0 = *(prefix_t *)(_h);
41-
hit = prefix[pr0];
42-
43-
if (hit) {
44-
45-
if (lookup32) {
46-
off = lookup32[pr0];
47-
l32 = _h[0];
48-
st= off;
49-
ed = off + hit - 1;
50-
while(st<=ed) {
51-
mi = (st+ed)/2;
52-
lmi = lookup32[mi];
53-
if(l32<lmi) {
54-
ed = mi - 1;
55-
} else if(l32==lmi) {
56-
// found
57-
goto addItem;
58-
} else {
59-
st = mi + 1;
60-
}
61-
}
62-
return;
39+
char add[48];
40+
41+
if (prefix == NULL) {
42+
43+
// No lookup compute address and return
44+
char *pattern = (char *)lookup32;
45+
_GetAddress(type, _h, add);
46+
if (_Match(add, pattern)) {
47+
// found
48+
goto addItem;
6349
}
50+
51+
} else {
52+
53+
// Lookup table
54+
pr0 = *(prefix_t *)(_h);
55+
hit = prefix[pr0];
56+
57+
if (hit) {
58+
59+
if (lookup32) {
60+
off = lookup32[pr0];
61+
l32 = _h[0];
62+
st = off;
63+
ed = off + hit - 1;
64+
while (st <= ed) {
65+
mi = (st + ed) / 2;
66+
lmi = lookup32[mi];
67+
if (l32 < lmi) {
68+
ed = mi - 1;
69+
} else if (l32 == lmi) {
70+
// found
71+
goto addItem;
72+
} else {
73+
st = mi + 1;
74+
}
75+
}
76+
return;
77+
}
78+
79+
addItem:
80+
81+
pos = atomicAdd(out, 1);
82+
if (pos < maxFound) {
83+
out[pos*ITEM_SIZE32 + 1] = tid;
84+
out[pos*ITEM_SIZE32 + 2] = (uint32_t)(incr << 16) | (uint32_t)(mode << 15) | (uint32_t)(endo);
85+
out[pos*ITEM_SIZE32 + 3] = _h[0];
86+
out[pos*ITEM_SIZE32 + 4] = _h[1];
87+
out[pos*ITEM_SIZE32 + 5] = _h[2];
88+
out[pos*ITEM_SIZE32 + 6] = _h[3];
89+
out[pos*ITEM_SIZE32 + 7] = _h[4];
90+
}
6491

65-
addItem:
66-
67-
pos = atomicAdd(out, 1);
68-
if (pos < maxFound) {
69-
out[pos*ITEM_SIZE32 + 1] = tid;
70-
out[pos*ITEM_SIZE32 + 2] = (uint32_t)(incr << 16) | (uint32_t)(mode << 15) | (uint32_t)(endo);
71-
out[pos*ITEM_SIZE32 + 3] = _h[0];
72-
out[pos*ITEM_SIZE32 + 4] = _h[1];
73-
out[pos*ITEM_SIZE32 + 5] = _h[2];
74-
out[pos*ITEM_SIZE32 + 6] = _h[3];
75-
out[pos*ITEM_SIZE32 + 7] = _h[4];
7692
}
7793

7894
}
@@ -81,7 +97,8 @@ __device__ __noinline__ void CheckPoint(uint32_t *_h, int32_t incr, int32_t endo
8197

8298
// -----------------------------------------------------------------------------------------
8399

84-
#define CHECK_POINT(_h,incr,endo,mode) CheckPoint(_h,incr,endo,mode,prefix,lookup32,maxFound,out)
100+
#define CHECK_POINT(_h,incr,endo,mode) CheckPoint(_h,incr,endo,mode,prefix,lookup32,maxFound,out,P2PKH)
101+
#define CHECK_POINT_P2SH(_h,incr,endo,mode) CheckPoint(_h,incr,endo,mode,prefix,lookup32,maxFound,out,P2SH)
85102

86103
__device__ __noinline__ void CheckHashComp(prefix_t *prefix, uint64_t *px, uint8_t isOdd, int32_t incr,
87104
uint32_t *lookup32, uint32_t maxFound, uint32_t *out) {
@@ -106,6 +123,7 @@ __device__ __noinline__ void CheckHashComp(prefix_t *prefix, uint64_t *px, uint8
106123
_GetHash160Comp(pe2x, !isOdd, (uint8_t *)h);
107124
CHECK_POINT(h, -incr, 2, true);
108125

126+
109127
}
110128

111129
__device__ __noinline__ void CheckHashP2SHComp(prefix_t *prefix, uint64_t *px, uint8_t isOdd, int32_t incr,
@@ -116,20 +134,20 @@ __device__ __noinline__ void CheckHashP2SHComp(prefix_t *prefix, uint64_t *px, u
116134
uint64_t pe2x[4];
117135

118136
_GetHash160P2SHComp(px, isOdd, (uint8_t *)h);
119-
CHECK_POINT(h, incr, 0, true);
137+
CHECK_POINT_P2SH(h, incr, 0, true);
120138
_ModMult(pe1x, px, _beta);
121139
_GetHash160P2SHComp(pe1x, isOdd, (uint8_t *)h);
122-
CHECK_POINT(h, incr, 1, true);
140+
CHECK_POINT_P2SH(h, incr, 1, true);
123141
_ModMult(pe2x, px, _beta2);
124142
_GetHash160P2SHComp(pe2x, isOdd, (uint8_t *)h);
125-
CHECK_POINT(h, incr, 2, true);
143+
CHECK_POINT_P2SH(h, incr, 2, true);
126144

127145
_GetHash160P2SHComp(px, !isOdd, (uint8_t *)h);
128-
CHECK_POINT(h, -incr, 0, true);
146+
CHECK_POINT_P2SH(h, -incr, 0, true);
129147
_GetHash160P2SHComp(pe1x, !isOdd, (uint8_t *)h);
130-
CHECK_POINT(h, -incr, 1, true);
148+
CHECK_POINT_P2SH(h, -incr, 1, true);
131149
_GetHash160P2SHComp(pe2x, !isOdd, (uint8_t *)h);
132-
CHECK_POINT(h, -incr, 2, true);
150+
CHECK_POINT_P2SH(h, -incr, 2, true);
133151

134152
}
135153

@@ -172,22 +190,22 @@ __device__ __noinline__ void CheckHashP2SHUncomp(prefix_t *prefix, uint64_t *px,
172190
uint64_t pyn[4];
173191

174192
_GetHash160P2SHUncomp(px, py, (uint8_t *)h);
175-
CHECK_POINT(h, incr, 0, false);
193+
CHECK_POINT_P2SH(h, incr, 0, false);
176194
_ModMult(pe1x, px, _beta);
177195
_GetHash160P2SHUncomp(pe1x, py, (uint8_t *)h);
178-
CHECK_POINT(h, incr, 1, false);
196+
CHECK_POINT_P2SH(h, incr, 1, false);
179197
_ModMult(pe2x, px, _beta2);
180198
_GetHash160P2SHUncomp(pe2x, py, (uint8_t *)h);
181-
CHECK_POINT(h, incr, 2, false);
199+
CHECK_POINT_P2SH(h, incr, 2, false);
182200

183201
ModNeg256(pyn, py);
184202

185203
_GetHash160P2SHUncomp(px, pyn, (uint8_t *)h);
186-
CHECK_POINT(h, -incr, 0, false);
204+
CHECK_POINT_P2SH(h, -incr, 0, false);
187205
_GetHash160P2SHUncomp(pe1x, pyn, (uint8_t *)h);
188-
CHECK_POINT(h, -incr, 1, false);
206+
CHECK_POINT_P2SH(h, -incr, 1, false);
189207
_GetHash160P2SHUncomp(pe2x, pyn, (uint8_t *)h);
190-
CHECK_POINT(h, -incr, 2, false);
208+
CHECK_POINT_P2SH(h, -incr, 2, false);
191209

192210
}
193211

@@ -245,6 +263,7 @@ __device__ void ComputeKeys(uint32_t mode, uint64_t *startx, uint64_t *starty,
245263
uint64_t dy[4];
246264
uint64_t _s[4];
247265
uint64_t _p2[4];
266+
char pattern[48];
248267

249268
// Load starting key
250269
__syncthreads();
@@ -253,6 +272,11 @@ __device__ void ComputeKeys(uint32_t mode, uint64_t *startx, uint64_t *starty,
253272
Load256(px, sx);
254273
Load256(py, sy);
255274

275+
if (sPrefix == NULL) {
276+
memcpy(pattern,lookup32,48);
277+
lookup32 = (uint32_t *)pattern;
278+
}
279+
256280
for (uint32_t j = 0; j < STEP_SIZE / GRP_SIZE; j++) {
257281

258282
// Fill group with delta x
@@ -370,6 +394,7 @@ __device__ void ComputeKeysP2SH(uint32_t mode, uint64_t *startx, uint64_t *start
370394
uint64_t dy[4];
371395
uint64_t _s[4];
372396
uint64_t _p2[4];
397+
char pattern[48];
373398

374399
// Load starting key
375400
__syncthreads();
@@ -378,6 +403,11 @@ __device__ void ComputeKeysP2SH(uint32_t mode, uint64_t *startx, uint64_t *start
378403
Load256(px, sx);
379404
Load256(py, sy);
380405

406+
if (sPrefix == NULL) {
407+
memcpy(pattern, lookup32, 48);
408+
lookup32 = (uint32_t *)pattern;
409+
}
410+
381411
for (uint32_t j = 0; j < STEP_SIZE / GRP_SIZE; j++) {
382412

383413
// Fill group with delta x

0 commit comments

Comments
 (0)
Please sign in to comment.