20
20
// For the kernel, we use a 16 bits prefix lookup table which correspond to ~3 Base58 characters
21
21
// A second level lookup table contains 32 bits prefix (if used)
22
22
// (The CPU computes the full address and check the full prefix)
23
- //
23
+ //
24
24
// We use affine coordinates for elliptic curve point (ie Z=1)
25
25
26
- __device__ __noinline__ void CheckPoint (uint32_t * _h , int32_t incr , int32_t endo , int32_t mode ,prefix_t * prefix ,
26
+ __device__ __noinline__ void CheckPoint (uint32_t * _h , int32_t incr , int32_t endo , int32_t mode ,prefix_t * prefix ,
27
27
uint32_t * lookup32 , uint32_t maxFound , uint32_t * out ,int type ) {
28
28
29
29
uint32_t off ;
@@ -37,7 +37,7 @@ __device__ __noinline__ void CheckPoint(uint32_t *_h, int32_t incr, int32_t endo
37
37
uint32_t lmi ;
38
38
uint32_t tid = (blockIdx .x * blockDim .x ) + threadIdx .x ;
39
39
char add [48 ];
40
-
40
+
41
41
if (prefix == NULL ) {
42
42
43
43
// No lookup compute address and return
@@ -47,9 +47,9 @@ __device__ __noinline__ void CheckPoint(uint32_t *_h, int32_t incr, int32_t endo
47
47
// found
48
48
goto addItem ;
49
49
}
50
-
50
+
51
51
} else {
52
-
52
+
53
53
// Lookup table
54
54
pr0 = * (prefix_t * )(_h );
55
55
hit = prefix [pr0 ];
@@ -100,7 +100,7 @@ __device__ __noinline__ void CheckPoint(uint32_t *_h, int32_t incr, int32_t endo
100
100
#define CHECK_POINT (_h ,incr ,endo ,mode ) CheckPoint(_h,incr,endo,mode,prefix,lookup32,maxFound,out,P2PKH)
101
101
#define CHECK_POINT_P2SH (_h ,incr ,endo ,mode ) CheckPoint(_h,incr,endo,mode,prefix,lookup32,maxFound,out,P2SH)
102
102
103
- __device__ __noinline__ void CheckHashComp (prefix_t * prefix , uint64_t * px , uint8_t isOdd , int32_t incr ,
103
+ __device__ __noinline__ void CheckHashComp (prefix_t * prefix , uint64_t * px , uint8_t isOdd , int32_t incr ,
104
104
uint32_t * lookup32 , uint32_t maxFound , uint32_t * out ) {
105
105
106
106
uint32_t h [5 ];
@@ -153,7 +153,7 @@ __device__ __noinline__ void CheckHashP2SHComp(prefix_t *prefix, uint64_t *px, u
153
153
154
154
// -----------------------------------------------------------------------------------------
155
155
156
- __device__ __noinline__ void CheckHashUncomp (prefix_t * prefix , uint64_t * px , uint64_t * py , int32_t incr ,
156
+ __device__ __noinline__ void CheckHashUncomp (prefix_t * prefix , uint64_t * px , uint64_t * py , int32_t incr ,
157
157
uint32_t * lookup32 , uint32_t maxFound , uint32_t * out ) {
158
158
159
159
uint32_t h [5 ];
@@ -211,7 +211,7 @@ __device__ __noinline__ void CheckHashP2SHUncomp(prefix_t *prefix, uint64_t *px,
211
211
212
212
// -----------------------------------------------------------------------------------------
213
213
214
- __device__ __noinline__ void CheckHash (uint32_t mode , prefix_t * prefix , uint64_t * px , uint64_t * py , int32_t incr ,
214
+ __device__ __noinline__ void CheckHash (uint32_t mode , prefix_t * prefix , uint64_t * px , uint64_t * py , int32_t incr ,
215
215
uint32_t * lookup32 , uint32_t maxFound , uint32_t * out ) {
216
216
217
217
switch (mode ) {
@@ -251,7 +251,7 @@ __device__ __noinline__ void CheckP2SHHash(uint32_t mode, prefix_t *prefix, uint
251
251
252
252
// -----------------------------------------------------------------------------------------
253
253
254
- __device__ void ComputeKeys (uint32_t mode , uint64_t * startx , uint64_t * starty ,
254
+ __device__ void ComputeKeys (uint32_t mode , uint64_t * startx , uint64_t * starty ,
255
255
prefix_t * sPrefix , uint32_t * lookup32 , uint32_t maxFound , uint32_t * out ) {
256
256
257
257
uint64_t dx [GRP_SIZE /2 + 1 ][4 ];
@@ -312,7 +312,7 @@ __device__ void ComputeKeys(uint32_t mode, uint64_t *startx, uint64_t *starty,
312
312
313
313
ModSub256 (py , Gx [i ], px );
314
314
_ModMult (py , _s ); // py = - s*(ret.x-p2.x)
315
- ModSub256 (py , Gy [i ]); // py = - p2.y - s*(ret.x-p2.x);
315
+ ModSub256 (py , Gy [i ]); // py = - p2.y - s*(ret.x-p2.x);
316
316
317
317
CHECK_PREFIX (GRP_SIZE / 2 + (i + 1 ));
318
318
@@ -328,7 +328,7 @@ __device__ void ComputeKeys(uint32_t mode, uint64_t *startx, uint64_t *starty,
328
328
329
329
ModSub256 (py , Gx [i ], px );
330
330
_ModMult (py , _s ); // py = - s*(ret.x-p2.x)
331
- ModAdd256 (py , Gy [i ]); // py = - p2.y - s*(ret.x-p2.x);
331
+ ModAdd256 (py , Gy [i ]); // py = - p2.y - s*(ret.x-p2.x);
332
332
333
333
CHECK_PREFIX (GRP_SIZE / 2 - (i + 1 ));
334
334
@@ -348,7 +348,7 @@ __device__ void ComputeKeys(uint32_t mode, uint64_t *startx, uint64_t *starty,
348
348
349
349
ModSub256 (py , Gx [i ], px );
350
350
_ModMult (py , _s ); // py = - s*(ret.x-p2.x)
351
- ModAdd256 (py , Gy [i ]); // py = - p2.y - s*(ret.x-p2.x);
351
+ ModAdd256 (py , Gy [i ]); // py = - p2.y - s*(ret.x-p2.x);
352
352
353
353
CHECK_PREFIX (0 );
354
354
@@ -367,7 +367,7 @@ __device__ void ComputeKeys(uint32_t mode, uint64_t *startx, uint64_t *starty,
367
367
368
368
ModSub256 (py , _2Gnx , px );
369
369
_ModMult (py , _s ); // py = - s*(ret.x-p2.x)
370
- ModSub256 (py , _2Gny ); // py = - p2.y - s*(ret.x-p2.x);
370
+ ModSub256 (py , _2Gny ); // py = - p2.y - s*(ret.x-p2.x);
371
371
372
372
}
373
373
@@ -443,7 +443,7 @@ __device__ void ComputeKeysP2SH(uint32_t mode, uint64_t *startx, uint64_t *start
443
443
444
444
ModSub256 (py , Gx [i ], px );
445
445
_ModMult (py , _s ); // py = - s*(ret.x-p2.x)
446
- ModSub256 (py , Gy [i ]); // py = - p2.y - s*(ret.x-p2.x);
446
+ ModSub256 (py , Gy [i ]); // py = - p2.y - s*(ret.x-p2.x);
447
447
448
448
CHECK_PREFIX_P2SH (GRP_SIZE / 2 + (i + 1 ));
449
449
@@ -459,7 +459,7 @@ __device__ void ComputeKeysP2SH(uint32_t mode, uint64_t *startx, uint64_t *start
459
459
460
460
ModSub256 (py , Gx [i ], px );
461
461
_ModMult (py , _s ); // py = - s*(ret.x-p2.x)
462
- ModAdd256 (py , Gy [i ]); // py = - p2.y - s*(ret.x-p2.x);
462
+ ModAdd256 (py , Gy [i ]); // py = - p2.y - s*(ret.x-p2.x);
463
463
464
464
CHECK_PREFIX_P2SH (GRP_SIZE / 2 - (i + 1 ));
465
465
@@ -479,7 +479,7 @@ __device__ void ComputeKeysP2SH(uint32_t mode, uint64_t *startx, uint64_t *start
479
479
480
480
ModSub256 (py , Gx [i ], px );
481
481
_ModMult (py , _s ); // py = - s*(ret.x-p2.x)
482
- ModAdd256 (py , Gy [i ]); // py = - p2.y - s*(ret.x-p2.x);
482
+ ModAdd256 (py , Gy [i ]); // py = - p2.y - s*(ret.x-p2.x);
483
483
484
484
CHECK_PREFIX_P2SH (0 );
485
485
@@ -498,7 +498,7 @@ __device__ void ComputeKeysP2SH(uint32_t mode, uint64_t *startx, uint64_t *start
498
498
499
499
ModSub256 (py , _2Gnx , px );
500
500
_ModMult (py , _s ); // py = - s*(ret.x-p2.x)
501
- ModSub256 (py , _2Gny ); // py = - p2.y - s*(ret.x-p2.x);
501
+ ModSub256 (py , _2Gny ); // py = - p2.y - s*(ret.x-p2.x);
502
502
503
503
}
504
504
@@ -629,7 +629,7 @@ __device__ void ComputeKeysComp(uint64_t *startx, uint64_t *starty, prefix_t *sP
629
629
630
630
ModSub256 (py , _2Gnx , px );
631
631
_ModMult (py , _s ); // py = - s*(ret.x-p2.x)
632
- ModSub256 (py , _2Gny ); // py = - p2.y - s*(ret.x-p2.x);
632
+ ModSub256 (py , _2Gny ); // py = - p2.y - s*(ret.x-p2.x);
633
633
634
634
}
635
635
0 commit comments