Skip to content

Commit f8aa16f

Browse files
tpruvottpruvot
committed
skein: cleanup, and precompute h8
1 parent 683dc0e commit f8aa16f

File tree

1 file changed

+69
-59
lines changed

1 file changed

+69
-59
lines changed

Diff for: ‎quark/cuda_skein512.cu

+69-59
Original file line numberDiff line numberDiff line change
@@ -417,9 +417,9 @@ void quark_skein512_gpu_hash_64(uint32_t threads, uint32_t startNounce, uint64_t
417417
uint32_t nounce = (g_nonceVector != NULL) ? g_nonceVector[thread] : (startNounce + thread);
418418

419419
uint32_t hashPosition = nounce - startNounce;
420-
uint64_t *inpHash = &g_hash[hashPosition * 8U];
420+
uint2 *inpHash = (uint2*) (&g_hash[hashPosition * 8U]);
421421

422-
// Initialisierung
422+
// Init
423423
h0 = vectorize(0x4903ADFF749C51CEull);
424424
h1 = vectorize(0x0D95DE399746DF03ull);
425425
h2 = vectorize(0x8FD1934127C79BCEull);
@@ -433,17 +433,19 @@ void quark_skein512_gpu_hash_64(uint32_t threads, uint32_t startNounce, uint64_t
433433
// 1st Round -> etype = 480, ptr = 64, bcount = 0, data = msg
434434
#pragma unroll 8
435435
for (int i = 0; i < 8; i++)
436-
p[i] = vectorize(inpHash[i]);
436+
p[i] = inpHash[i];
437437

438-
t0 = vectorize(64); // ptr
439-
// t1 = vectorize(480ull << 55); // etype
440-
t1 = vectorize(0xf000000000000000ULL);
438+
t0 = make_uint2(0x40, 0); // 64
439+
t1 = vectorize(0xf000000000000000ULL); // 480ull << 55 (etype)
441440

442441
//#if CUDA_VERSION >= 7000
443442
// doesnt really affect x11 perfs.
444443
__threadfence();
445444
//#endif
446-
TFBIG_KINIT_UI2(h0, h1, h2, h3, h4, h5, h6, h7, h8, t0, t1, t2);
445+
//TFBIG_KINIT_UI2(h0, h1, h2, h3, h4, h5, h6, h7, h8, t0, t1, t2);
446+
h8 = vectorize(0xcab2076d98173ec4ULL);
447+
t2 = vectorize(0xf000000000000040ULL);
448+
447449
TFBIG_4e_UI2(0);
448450
TFBIG_4o_UI2(1);
449451
TFBIG_4e_UI2(2);
@@ -464,23 +466,22 @@ void quark_skein512_gpu_hash_64(uint32_t threads, uint32_t startNounce, uint64_t
464466
TFBIG_4o_UI2(17);
465467
TFBIG_ADDKEY_UI2(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, 18);
466468

467-
h0 = vectorize(inpHash[0]) ^ p[0];
468-
h1 = vectorize(inpHash[1]) ^ p[1];
469-
h2 = vectorize(inpHash[2]) ^ p[2];
470-
h3 = vectorize(inpHash[3]) ^ p[3];
471-
h4 = vectorize(inpHash[4]) ^ p[4];
472-
h5 = vectorize(inpHash[5]) ^ p[5];
473-
h6 = vectorize(inpHash[6]) ^ p[6];
474-
h7 = vectorize(inpHash[7]) ^ p[7];
469+
h0 = inpHash[0] ^ p[0];
470+
h1 = inpHash[1] ^ p[1];
471+
h2 = inpHash[2] ^ p[2];
472+
h3 = inpHash[3] ^ p[3];
473+
h4 = inpHash[4] ^ p[4];
474+
h5 = inpHash[5] ^ p[5];
475+
h6 = inpHash[6] ^ p[6];
476+
h7 = inpHash[7] ^ p[7];
475477

476478
// 2. Runde -> etype = 510, ptr = 8, bcount = 0, data = 0
477479
#pragma unroll 8
478480
for(int i=0; i<8; i++)
479481
p[i] = vectorize(0);
480482

481-
t0 = vectorize(8); // ptr
482-
//t1 = vectorize(510ull << 55); // etype
483-
t1 = vectorize(0xff00000000000000ULL);
483+
t0 = make_uint2(0x8, 0);
484+
t1 = vectorize(0xff00000000000000ULL); // etype
484485

485486
TFBIG_KINIT_UI2(h0, h1, h2, h3, h4, h5, h6, h7, h8, t0, t1, t2);
486487
TFBIG_4e_UI2(0);
@@ -526,7 +527,7 @@ void quark_skein512_gpu_hash_64_sm3(uint32_t threads, uint32_t startNounce, uint
526527
uint32_t nounce = (g_nonceVector != NULL) ? g_nonceVector[thread] : (startNounce + thread);
527528

528529
uint32_t hashPosition = nounce - startNounce;
529-
uint64_t *inpHash = &g_hash[hashPosition * 8];
530+
uint64_t *inpHash = &g_hash[hashPosition * 8U];
530531

531532
// Init
532533
h0 = 0x4903ADFF749C51CEull;
@@ -538,16 +539,18 @@ void quark_skein512_gpu_hash_64_sm3(uint32_t threads, uint32_t startNounce, uint
538539
h6 = 0x991112C71A75B523ull;
539540
h7 = 0xAE18A40B660FCC33ull;
540541

541-
// 1. Runde -> etype = 480, ptr = 64, bcount = 0, data = msg
542+
// 1st Round -> etype = 480, ptr = 64, bcount = 0, data = msg
542543
#pragma unroll 8
543544
for(int i=0; i<8; i++)
544545
p[i] = inpHash[i];
545546

546-
t0 = 64; // ptr
547-
// t1 = 480ull << 55; // etype
548-
t1 = 0xf000000000000000ULL;
547+
t0 = 0x40; // 64.
548+
t1 = 0xf000000000000000ULL; // 480ull << 55 (etype)
549+
550+
//TFBIG_KINIT(h0, h1, h2, h3, h4, h5, h6, h7, h8, t0, t1, t2);
551+
h8 = 0xcab2076d98173ec4ULL;
552+
t2 = 0xf000000000000040ULL;
549553

550-
TFBIG_KINIT(h0, h1, h2, h3, h4, h5, h6, h7, h8, t0, t1, t2);
551554
TFBIG_4e(0);
552555
TFBIG_4o(1);
553556
TFBIG_4e(2);
@@ -577,14 +580,15 @@ void quark_skein512_gpu_hash_64_sm3(uint32_t threads, uint32_t startNounce, uint
577580
h6 = inpHash[6] ^ p[6];
578581
h7 = inpHash[7] ^ p[7];
579582

580-
// 2. Runde -> etype = 510, ptr = 8, bcount = 0, data = 0
583+
// 2nd Round -> etype = 510, ptr = 8, bcount = 0, data = 0
581584
#pragma unroll 8
582585
for(int i=0; i<8; i++)
583586
p[i] = 0ull;
584587

585588
t0 = 8; // ptr
586589
t1 = 510ull << 55; // etype
587590
TFBIG_KINIT(h0, h1, h2, h3, h4, h5, h6, h7, h8, t0, t1, t2);
591+
588592
TFBIG_4e(0);
589593
TFBIG_4o(1);
590594
TFBIG_4e(2);
@@ -606,7 +610,7 @@ void quark_skein512_gpu_hash_64_sm3(uint32_t threads, uint32_t startNounce, uint
606610
TFBIG_ADDKEY(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, 18);
607611

608612
// output
609-
uint64_t *outpHash = &g_hash[hashPosition * 8];
613+
uint64_t *outpHash = &g_hash[hashPosition * 8U];
610614

611615
#pragma unroll 8
612616
for(int i=0; i<8; i++)
@@ -633,8 +637,6 @@ void skein512_gpu_hash_80(uint32_t threads, uint32_t startNounce, uint64_t *outp
633637
h6 = vectorize(c_PaddedMessage80[16]);
634638
h7 = vectorize(c_PaddedMessage80[17]);
635639

636-
t2 = vectorize(c_PaddedMessage80[18]);
637-
638640
uint32_t nonce = swap ? cuda_swab32(startNounce + thread) : startNounce + thread;
639641
uint2 nonce2 = make_uint2(_LODWORD(c_PaddedMessage80[9]), nonce);
640642

@@ -646,9 +648,13 @@ void skein512_gpu_hash_80(uint32_t threads, uint32_t startNounce, uint64_t *outp
646648
for (int i = 2; i < 8; i++)
647649
p[i] = vectorize(0ull);
648650

649-
t0 = vectorize(0x50ull);
651+
t0 = make_uint2(0x50, 0);
650652
t1 = vectorize(0xB000000000000000ull);
651-
TFBIG_KINIT_UI2(h0, h1, h2, h3, h4, h5, h6, h7, h8, t0, t1, t2);
653+
654+
//TFBIG_KINIT_UI2(h0, h1, h2, h3, h4, h5, h6, h7, h8, t0, t1, t2);
655+
h8 = vectorize(c_PaddedMessage80[18]);
656+
t2 = vectorize(0xB000000000000050ull); // t0 ^ t1
657+
652658
TFBIG_4e_UI2(0);
653659
TFBIG_4o_UI2(1);
654660
TFBIG_4e_UI2(2);
@@ -669,7 +675,7 @@ void skein512_gpu_hash_80(uint32_t threads, uint32_t startNounce, uint64_t *outp
669675
TFBIG_4o_UI2(17);
670676
TFBIG_ADDKEY_UI2(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, 18);
671677

672-
uint64_t *outpHash = &output64[thread * 8];
678+
uint64_t *outpHash = &output64[thread * 8U];
673679
outpHash[0] = c_PaddedMessage80[8] ^ devectorize(p[0]);
674680
outpHash[1] = devectorize(nonce2 ^ p[1]);
675681
#pragma unroll
@@ -684,23 +690,22 @@ void skein512_gpu_hash_80_sm3(uint32_t threads, uint32_t startNounce, uint64_t *
684690
uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x);
685691
if (thread < threads)
686692
{
687-
uint64_t h0, h1, h2, h3, h4, h5, h6, h7, h8;
688-
uint64_t t0, t1, t2;
689-
690693
// Init
691-
h0 = 0x4903ADFF749C51CEull;
692-
h1 = 0x0D95DE399746DF03ull;
693-
h2 = 0x8FD1934127C79BCEull;
694-
h3 = 0x9A255629FF352CB1ull;
695-
h4 = 0x5DB62599DF6CA7B0ull;
696-
h5 = 0xEABE394CA9D5C3F4ull;
697-
h6 = 0x991112C71A75B523ull;
698-
h7 = 0xAE18A40B660FCC33ull;
699-
700-
t0 = 64; // ptr
701-
//t1 = vectorize(0xE0ull << 55); // etype
702-
t1 = 0x7000000000000000ull;
703-
TFBIG_KINIT(h0, h1, h2, h3, h4, h5, h6, h7, h8, t0, t1, t2);
694+
uint64_t h0 = 0x4903ADFF749C51CEull;
695+
uint64_t h1 = 0x0D95DE399746DF03ull;
696+
uint64_t h2 = 0x8FD1934127C79BCEull;
697+
uint64_t h3 = 0x9A255629FF352CB1ull;
698+
uint64_t h4 = 0x5DB62599DF6CA7B0ull;
699+
uint64_t h5 = 0xEABE394CA9D5C3F4ull;
700+
uint64_t h6 = 0x991112C71A75B523ull;
701+
uint64_t h7 = 0xAE18A40B660FCC33ull;
702+
703+
uint64_t t0 = 0x40; // ptr = 64.
704+
uint64_t t1 = 0x7000000000000000ull; // 0xE0ull << 55 // etype
705+
706+
//TFBIG_KINIT(h0, h1, h2, h3, h4, h5, h6, h7, h8, t0, t1, t2);
707+
uint64_t t2 = 0x7000000000000040ull;
708+
uint64_t h8 = 0xcab2076d98173ec4ull;
704709

705710
uint64_t p[8];
706711
#pragma unroll 8
@@ -745,12 +750,15 @@ void skein512_gpu_hash_80_sm3(uint32_t threads, uint32_t startNounce, uint64_t *
745750

746751
#pragma unroll
747752
for (int i = 2; i < 8; i++)
748-
p[i] = 0ull;
753+
p[i] = 0;
749754

750-
t0 = 0x50ull; // SPH_T64(bcount << 6) + (sph_u64)(extra);
755+
t0 = 0x50; // (bcount << 6) + extra;
751756
t1 = 0xB000000000000000ull; // (bcount >> 58) + ((sph_u64)(etype) << 55);
752757

753-
TFBIG_KINIT(h0, h1, h2, h3, h4, h5, h6, h7, h8, t0, t1, t2);
758+
//TFBIG_KINIT(h0, h1, h2, h3, h4, h5, h6, h7, h8, t0, t1, t2);
759+
t2 = 0xB000000000000050ull;
760+
h8 = c_PaddedMessage80[18];
761+
754762
TFBIG_4e(0);
755763
TFBIG_4o(1);
756764
TFBIG_4e(2);
@@ -773,7 +781,7 @@ void skein512_gpu_hash_80_sm3(uint32_t threads, uint32_t startNounce, uint64_t *
773781

774782
// skein_big_close 2nd loop -> etype = 0x1fe, ptr = 8, bcount = 0
775783
// output
776-
uint64_t *outpHash = &output64[thread * 8];
784+
uint64_t *outpHash = &output64[thread * 8U];
777785
outpHash[0] = c_PaddedMessage80[8] ^ p[0];
778786
outpHash[1] = nonce64 ^ p[1];
779787
#pragma unroll
@@ -788,11 +796,10 @@ void skein512_gpu_hash_close(uint32_t threads, uint32_t startNounce, uint64_t *g
788796
uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x);
789797
if (thread < threads)
790798
{
791-
uint2 t0 = vectorize(8); // extra
799+
uint2 t0 = make_uint2(0x8, 0); // extra
792800
uint2 t1 = vectorize(0xFF00000000000000ull); // etype
793-
uint2 t2 = vectorize(0xB000000000000050ull);
794801

795-
uint64_t *state = &g_hash[thread * 8];
802+
uint64_t *state = &g_hash[thread * 8U];
796803
uint2 h0 = vectorize(state[0]);
797804
uint2 h1 = vectorize(state[1]);
798805
uint2 h2 = vectorize(state[2]);
@@ -801,7 +808,8 @@ void skein512_gpu_hash_close(uint32_t threads, uint32_t startNounce, uint64_t *g
801808
uint2 h5 = vectorize(state[5]);
802809
uint2 h6 = vectorize(state[6]);
803810
uint2 h7 = vectorize(state[7]);
804-
uint2 h8;
811+
812+
uint2 h8, t2;
805813
TFBIG_KINIT_UI2(h0, h1, h2, h3, h4, h5, h6, h7, h8, t0, t1, t2);
806814

807815
uint2 p[8] = { 0 };
@@ -841,9 +849,8 @@ void skein512_gpu_hash_close_sm3(uint32_t threads, uint32_t startNounce, uint64_
841849
{
842850
uint64_t t0 = 8ull; // extra
843851
uint64_t t1 = 0xFF00000000000000ull; // etype
844-
uint64_t t2 = 0xB000000000000050ull;
845852

846-
uint64_t *state = &g_hash[thread * 8];
853+
uint64_t *state = &g_hash[thread * 8U];
847854

848855
uint64_t h0 = state[0];
849856
uint64_t h1 = state[1];
@@ -853,7 +860,7 @@ void skein512_gpu_hash_close_sm3(uint32_t threads, uint32_t startNounce, uint64_
853860
uint64_t h5 = state[5];
854861
uint64_t h6 = state[6];
855862
uint64_t h7 = state[7];
856-
uint64_t h8;
863+
uint64_t h8, t2;
857864
TFBIG_KINIT(h0, h1, h2, h3, h4, h5, h6, h7, h8, t0, t1, t2);
858865

859866
uint64_t p[8] = { 0 };
@@ -971,7 +978,10 @@ static void skein512_precalc_80(uint64_t* message)
971978
message[16] = message[6] ^ p[6];
972979
message[17] = message[7] ^ p[7];
973980

974-
message[18] = t2;
981+
// h8
982+
message[18] = 0x1BD11BDAA9FC1A22ULL;
983+
for (int i=10; i<18; i++)
984+
message[18] ^= message[i];
975985
}
976986

977987
__host__

0 commit comments

Comments
 (0)