Skip to content

Commit

Permalink
upgrade blake routine in neoscrypt
Browse files Browse the repository at this point in the history
  • Loading branch information
@djm34
djm34 committed May 17, 2015
1 parent 821ea07 commit 817bcae
Showing 1 changed file with 156 additions and 32 deletions.
188 changes: 156 additions & 32 deletions neoscrypt/cuda_neoscrypt.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -89,6 +89,26 @@ idx = BLAKE2S_SIGMA[idx0][idx1+1]; a += key[idx]; \
}
#endif

#if __CUDA_ARCH__ >= 500
#define BLAKE_G_PRE(idx0,idx1, a, b, c, d, key) { \
a += key[idx0]; \
a += b; d = __byte_perm(d^a,0, 0x1032); \
c += d; b = rotateR(b^c, 12); \
a += key[idx1]; \
a += b; d = __byte_perm(d^a,0, 0x0321); \
c += d; b = rotateR(b^c, 7); \
}
#else
#define BLAKE_G_PRE(idx0, idx1, a, b, c, d, key) { \
a += key[idx0]; \
a += b; d = rotate(d^a,16); \
c += d; b = rotateR(b^c, 12); \
a += key[idx1]; \
a += b; d = rotateR(d^a,8); \
c += d; b = rotateR(b^c, 7); \
}
#endif

#define ROTL32(x, n) ((x) << (n)) | ((x) >> (32 - (n)))
#define ROTR32(x, n) (((x) >> (n)) | ((x) << (32 - (n))))

Expand All @@ -105,41 +125,147 @@ idx = BLAKE2S_SIGMA_host[idx0][idx1+1]; a += key[idx]; \
static __forceinline__ __device__ void Blake2S(uint32_t * inout, const uint32_t * TheKey)
{
uint16 V;
uint32_t idx;
uint32_t idx;
uint8 tmpblock;

V.hi = BLAKE2S_IV_Vec;
V.lo = BLAKE2S_IV_Vec;
V.hi = BLAKE2S_IV_Vec;
V.lo = BLAKE2S_IV_Vec;
V.lo.s0 ^= 0x01012020;

// Copy input block for later
tmpblock = V.lo;

V.hi.s4 ^= BLAKE2S_BLOCK_SIZE;

for (int x = 0; x < 10; ++x)
{
BLAKE_G(x, 0x00, V.lo.s0, V.lo.s4, V.hi.s0, V.hi.s4, TheKey);
BLAKE_G(x, 0x02, V.lo.s1, V.lo.s5, V.hi.s1, V.hi.s5, TheKey);
BLAKE_G(x, 0x04, V.lo.s2, V.lo.s6, V.hi.s2, V.hi.s6, TheKey);
BLAKE_G(x, 0x06, V.lo.s3, V.lo.s7, V.hi.s3, V.hi.s7, TheKey);
BLAKE_G(x, 0x08, V.lo.s0, V.lo.s5, V.hi.s2, V.hi.s7, TheKey);
BLAKE_G(x, 0x0A, V.lo.s1, V.lo.s6, V.hi.s3, V.hi.s4, TheKey);
BLAKE_G(x, 0x0C, V.lo.s2, V.lo.s7, V.hi.s0, V.hi.s5, TheKey);
BLAKE_G(x, 0x0E, V.lo.s3, V.lo.s4, V.hi.s1, V.hi.s6, TheKey);
}

// { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 },
BLAKE_G_PRE(0, 1, V.lo.s0, V.lo.s4, V.hi.s0, V.hi.s4, TheKey);
BLAKE_G_PRE(2, 3, V.lo.s1, V.lo.s5, V.hi.s1, V.hi.s5, TheKey);
BLAKE_G_PRE(4, 5, V.lo.s2, V.lo.s6, V.hi.s2, V.hi.s6, TheKey);
BLAKE_G_PRE(6, 7, V.lo.s3, V.lo.s7, V.hi.s3, V.hi.s7, TheKey);
BLAKE_G_PRE(8, 9, V.lo.s0, V.lo.s5, V.hi.s2, V.hi.s7, TheKey);
BLAKE_G_PRE(10, 11, V.lo.s1, V.lo.s6, V.hi.s3, V.hi.s4, TheKey);
BLAKE_G_PRE(12, 13, V.lo.s2, V.lo.s7, V.hi.s0, V.hi.s5, TheKey);
BLAKE_G_PRE(14, 15, V.lo.s3, V.lo.s4, V.hi.s1, V.hi.s6, TheKey);
// { 14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 },
BLAKE_G_PRE(14, 10, V.lo.s0, V.lo.s4, V.hi.s0, V.hi.s4, TheKey);
BLAKE_G_PRE(4, 8, V.lo.s1, V.lo.s5, V.hi.s1, V.hi.s5, TheKey);
BLAKE_G_PRE(9, 15, V.lo.s2, V.lo.s6, V.hi.s2, V.hi.s6, TheKey);
BLAKE_G_PRE(13, 6, V.lo.s3, V.lo.s7, V.hi.s3, V.hi.s7, TheKey);
BLAKE_G_PRE(1, 12, V.lo.s0, V.lo.s5, V.hi.s2, V.hi.s7, TheKey);
BLAKE_G_PRE(0, 2, V.lo.s1, V.lo.s6, V.hi.s3, V.hi.s4, TheKey);
BLAKE_G_PRE(11, 7, V.lo.s2, V.lo.s7, V.hi.s0, V.hi.s5, TheKey);
BLAKE_G_PRE(5, 3, V.lo.s3, V.lo.s4, V.hi.s1, V.hi.s6, TheKey);
// { 11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4 },
BLAKE_G_PRE(11, 8, V.lo.s0, V.lo.s4, V.hi.s0, V.hi.s4, TheKey);
BLAKE_G_PRE(12, 0, V.lo.s1, V.lo.s5, V.hi.s1, V.hi.s5, TheKey);
BLAKE_G_PRE(5, 2, V.lo.s2, V.lo.s6, V.hi.s2, V.hi.s6, TheKey);
BLAKE_G_PRE(15, 13, V.lo.s3, V.lo.s7, V.hi.s3, V.hi.s7, TheKey);
BLAKE_G_PRE(10, 14, V.lo.s0, V.lo.s5, V.hi.s2, V.hi.s7, TheKey);
BLAKE_G_PRE(3, 6, V.lo.s1, V.lo.s6, V.hi.s3, V.hi.s4, TheKey);
BLAKE_G_PRE(7, 1, V.lo.s2, V.lo.s7, V.hi.s0, V.hi.s5, TheKey);
BLAKE_G_PRE(9, 4, V.lo.s3, V.lo.s4, V.hi.s1, V.hi.s6, TheKey);
// { 7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8 },
BLAKE_G_PRE(7, 9, V.lo.s0, V.lo.s4, V.hi.s0, V.hi.s4, TheKey);
BLAKE_G_PRE(3, 1, V.lo.s1, V.lo.s5, V.hi.s1, V.hi.s5, TheKey);
BLAKE_G_PRE(13, 12, V.lo.s2, V.lo.s6, V.hi.s2, V.hi.s6, TheKey);
BLAKE_G_PRE(11, 14, V.lo.s3, V.lo.s7, V.hi.s3, V.hi.s7, TheKey);
BLAKE_G_PRE(2, 6, V.lo.s0, V.lo.s5, V.hi.s2, V.hi.s7, TheKey);
BLAKE_G_PRE(5, 10, V.lo.s1, V.lo.s6, V.hi.s3, V.hi.s4, TheKey);
BLAKE_G_PRE(4, 0, V.lo.s2, V.lo.s7, V.hi.s0, V.hi.s5, TheKey);
BLAKE_G_PRE(15, 8, V.lo.s3, V.lo.s4, V.hi.s1, V.hi.s6, TheKey);
// { 9, 0, 5, 7, 2, 4, 10, 15, 14, 1, 11, 12, 6, 8, 3, 13 },
BLAKE_G_PRE(9, 0, V.lo.s0, V.lo.s4, V.hi.s0, V.hi.s4, TheKey);
BLAKE_G_PRE(5, 7, V.lo.s1, V.lo.s5, V.hi.s1, V.hi.s5, TheKey);
BLAKE_G_PRE(2, 4, V.lo.s2, V.lo.s6, V.hi.s2, V.hi.s6, TheKey);
BLAKE_G_PRE(10, 15, V.lo.s3, V.lo.s7, V.hi.s3, V.hi.s7, TheKey);
BLAKE_G_PRE(14, 1, V.lo.s0, V.lo.s5, V.hi.s2, V.hi.s7, TheKey);
BLAKE_G_PRE(11, 12, V.lo.s1, V.lo.s6, V.hi.s3, V.hi.s4, TheKey);
BLAKE_G_PRE(6, 8, V.lo.s2, V.lo.s7, V.hi.s0, V.hi.s5, TheKey);
BLAKE_G_PRE(3, 13, V.lo.s3, V.lo.s4, V.hi.s1, V.hi.s6, TheKey);
// { 2, 12, 6, 10, 0, 11, 8, 3, 4, 13, 7, 5, 15, 14, 1, 9 },
BLAKE_G_PRE(2, 12, V.lo.s0, V.lo.s4, V.hi.s0, V.hi.s4, TheKey);
BLAKE_G_PRE(6, 10, V.lo.s1, V.lo.s5, V.hi.s1, V.hi.s5, TheKey);
BLAKE_G_PRE(0, 11, V.lo.s2, V.lo.s6, V.hi.s2, V.hi.s6, TheKey);
BLAKE_G_PRE(8, 3, V.lo.s3, V.lo.s7, V.hi.s3, V.hi.s7, TheKey);
BLAKE_G_PRE(4, 13, V.lo.s0, V.lo.s5, V.hi.s2, V.hi.s7, TheKey);
BLAKE_G_PRE(7, 5, V.lo.s1, V.lo.s6, V.hi.s3, V.hi.s4, TheKey);
BLAKE_G_PRE(15, 14, V.lo.s2, V.lo.s7, V.hi.s0, V.hi.s5, TheKey);
BLAKE_G_PRE(1, 9, V.lo.s3, V.lo.s4, V.hi.s1, V.hi.s6, TheKey);
// { 12, 5, 1, 15, 14, 13, 4, 10, 0, 7, 6, 3, 9, 2, 8, 11 },
BLAKE_G_PRE(12, 5, V.lo.s0, V.lo.s4, V.hi.s0, V.hi.s4, TheKey);
BLAKE_G_PRE(1, 15, V.lo.s1, V.lo.s5, V.hi.s1, V.hi.s5, TheKey);
BLAKE_G_PRE(14, 13, V.lo.s2, V.lo.s6, V.hi.s2, V.hi.s6, TheKey);
BLAKE_G_PRE(4, 10, V.lo.s3, V.lo.s7, V.hi.s3, V.hi.s7, TheKey);
BLAKE_G_PRE(0, 7, V.lo.s0, V.lo.s5, V.hi.s2, V.hi.s7, TheKey);
BLAKE_G_PRE(6, 3, V.lo.s1, V.lo.s6, V.hi.s3, V.hi.s4, TheKey);
BLAKE_G_PRE(9, 2, V.lo.s2, V.lo.s7, V.hi.s0, V.hi.s5, TheKey);
BLAKE_G_PRE(8, 11, V.lo.s3, V.lo.s4, V.hi.s1, V.hi.s6, TheKey);
// { 13, 11, 7, 14, 12, 1, 3, 9, 5, 0, 15, 4, 8, 6, 2, 10 },
BLAKE_G_PRE(13, 11, V.lo.s0, V.lo.s4, V.hi.s0, V.hi.s4, TheKey);
BLAKE_G_PRE(7, 14, V.lo.s1, V.lo.s5, V.hi.s1, V.hi.s5, TheKey);
BLAKE_G_PRE(12, 1, V.lo.s2, V.lo.s6, V.hi.s2, V.hi.s6, TheKey);
BLAKE_G_PRE(3, 9, V.lo.s3, V.lo.s7, V.hi.s3, V.hi.s7, TheKey);
BLAKE_G_PRE(5, 0, V.lo.s0, V.lo.s5, V.hi.s2, V.hi.s7, TheKey);
BLAKE_G_PRE(15, 4, V.lo.s1, V.lo.s6, V.hi.s3, V.hi.s4, TheKey);
BLAKE_G_PRE(8, 6, V.lo.s2, V.lo.s7, V.hi.s0, V.hi.s5, TheKey);
BLAKE_G_PRE(2, 10, V.lo.s3, V.lo.s4, V.hi.s1, V.hi.s6, TheKey);
// { 6, 15, 14, 9, 11, 3, 0, 8, 12, 2, 13, 7, 1, 4, 10, 5 },
BLAKE_G_PRE(6, 15, V.lo.s0, V.lo.s4, V.hi.s0, V.hi.s4, TheKey);
BLAKE_G_PRE(14, 9, V.lo.s1, V.lo.s5, V.hi.s1, V.hi.s5, TheKey);
BLAKE_G_PRE(11, 3, V.lo.s2, V.lo.s6, V.hi.s2, V.hi.s6, TheKey);
BLAKE_G_PRE(0, 8, V.lo.s3, V.lo.s7, V.hi.s3, V.hi.s7, TheKey);
BLAKE_G_PRE(12, 2, V.lo.s0, V.lo.s5, V.hi.s2, V.hi.s7, TheKey);
BLAKE_G_PRE(13, 7, V.lo.s1, V.lo.s6, V.hi.s3, V.hi.s4, TheKey);
BLAKE_G_PRE(1, 4, V.lo.s2, V.lo.s7, V.hi.s0, V.hi.s5, TheKey);
BLAKE_G_PRE(10, 5, V.lo.s3, V.lo.s4, V.hi.s1, V.hi.s6, TheKey);
// { 10, 2, 8, 4, 7, 6, 1, 5, 15, 11, 9, 14, 3, 12, 13, 0 },
BLAKE_G_PRE(10, 2, V.lo.s0, V.lo.s4, V.hi.s0, V.hi.s4, TheKey);
BLAKE_G_PRE(8, 4, V.lo.s1, V.lo.s5, V.hi.s1, V.hi.s5, TheKey);
BLAKE_G_PRE(7, 6, V.lo.s2, V.lo.s6, V.hi.s2, V.hi.s6, TheKey);
BLAKE_G_PRE(1, 5, V.lo.s3, V.lo.s7, V.hi.s3, V.hi.s7, TheKey);
BLAKE_G_PRE(15, 11, V.lo.s0, V.lo.s5, V.hi.s2, V.hi.s7, TheKey);
BLAKE_G_PRE(9, 14, V.lo.s1, V.lo.s6, V.hi.s3, V.hi.s4, TheKey);
BLAKE_G_PRE(3, 12, V.lo.s2, V.lo.s7, V.hi.s0, V.hi.s5, TheKey);
BLAKE_G_PRE(13, 0, V.lo.s3, V.lo.s4, V.hi.s1, V.hi.s6, TheKey);
V.lo ^= V.hi;
V.lo ^= tmpblock;


V.hi = BLAKE2S_IV_Vec;
tmpblock = V.lo;

V.hi.s4 ^= 128;
V.hi.s6 = ~V.hi.s6;

for (int x = 0; x < 10; ++x)
// { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 },
BLAKE_G_PRE(0, 1, V.lo.s0, V.lo.s4, V.hi.s0, V.hi.s4, inout);
BLAKE_G_PRE(2, 3, V.lo.s1, V.lo.s5, V.hi.s1, V.hi.s5, inout);
BLAKE_G_PRE(4, 5, V.lo.s2, V.lo.s6, V.hi.s2, V.hi.s6, inout);
BLAKE_G_PRE(6, 7, V.lo.s3, V.lo.s7, V.hi.s3, V.hi.s7, inout);
BLAKE_G_PRE(8, 9, V.lo.s0, V.lo.s5, V.hi.s2, V.hi.s7, inout);
BLAKE_G_PRE(10, 11, V.lo.s1, V.lo.s6, V.hi.s3, V.hi.s4, inout);
BLAKE_G_PRE(12, 13, V.lo.s2, V.lo.s7, V.hi.s0, V.hi.s5, inout);
BLAKE_G_PRE(14, 15, V.lo.s3, V.lo.s4, V.hi.s1, V.hi.s6, inout);
// { 14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 },
BLAKE_G_PRE(14, 10, V.lo.s0, V.lo.s4, V.hi.s0, V.hi.s4, inout);
BLAKE_G_PRE(4, 8, V.lo.s1, V.lo.s5, V.hi.s1, V.hi.s5, inout);
BLAKE_G_PRE(9, 15, V.lo.s2, V.lo.s6, V.hi.s2, V.hi.s6, inout);
BLAKE_G_PRE(13, 6, V.lo.s3, V.lo.s7, V.hi.s3, V.hi.s7, inout);
BLAKE_G_PRE(1, 12, V.lo.s0, V.lo.s5, V.hi.s2, V.hi.s7, inout);
BLAKE_G_PRE(0, 2, V.lo.s1, V.lo.s6, V.hi.s3, V.hi.s4, inout);
BLAKE_G_PRE(11, 7, V.lo.s2, V.lo.s7, V.hi.s0, V.hi.s5, inout);
BLAKE_G_PRE(5, 3, V.lo.s3, V.lo.s4, V.hi.s1, V.hi.s6, inout);
// { 11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4 },
BLAKE_G_PRE(11, 8, V.lo.s0, V.lo.s4, V.hi.s0, V.hi.s4, inout);
BLAKE_G_PRE(12, 0, V.lo.s1, V.lo.s5, V.hi.s1, V.hi.s5, inout);
BLAKE_G_PRE(5, 2, V.lo.s2, V.lo.s6, V.hi.s2, V.hi.s6, inout);
BLAKE_G_PRE(15, 13, V.lo.s3, V.lo.s7, V.hi.s3, V.hi.s7, inout);
BLAKE_G_PRE(10, 14, V.lo.s0, V.lo.s5, V.hi.s2, V.hi.s7, inout);
BLAKE_G_PRE(3, 6, V.lo.s1, V.lo.s6, V.hi.s3, V.hi.s4, inout);
BLAKE_G_PRE(7, 1, V.lo.s2, V.lo.s7, V.hi.s0, V.hi.s5, inout);
BLAKE_G_PRE(9, 4, V.lo.s3, V.lo.s4, V.hi.s1, V.hi.s6, inout);
// { 7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8 },
BLAKE_G_PRE(7, 9, V.lo.s0, V.lo.s4, V.hi.s0, V.hi.s4, inout);
BLAKE_G_PRE(3, 1, V.lo.s1, V.lo.s5, V.hi.s1, V.hi.s5, inout);
BLAKE_G_PRE(13, 12, V.lo.s2, V.lo.s6, V.hi.s2, V.hi.s6, inout);
BLAKE_G_PRE(11, 14, V.lo.s3, V.lo.s7, V.hi.s3, V.hi.s7, inout);
BLAKE_G_PRE(2, 6, V.lo.s0, V.lo.s5, V.hi.s2, V.hi.s7, inout);
BLAKE_G_PRE(5, 10, V.lo.s1, V.lo.s6, V.hi.s3, V.hi.s4, inout);
BLAKE_G_PRE(4, 0, V.lo.s2, V.lo.s7, V.hi.s0, V.hi.s5, inout);
BLAKE_G_PRE(15, 8, V.lo.s3, V.lo.s4, V.hi.s1, V.hi.s6, inout);
for (int x = 4; x < 10; ++x)
{
BLAKE_G(x, 0x00, V.lo.s0, V.lo.s4, V.hi.s0, V.hi.s4, inout);
BLAKE_G(x, 0x02, V.lo.s1, V.lo.s5, V.hi.s1, V.hi.s5, inout);
Expand All @@ -150,13 +276,11 @@ static __forceinline__ __device__ void Blake2S(uint32_t * inout, const uint32_t
BLAKE_G(x, 0x0C, V.lo.s2, V.lo.s7, V.hi.s0, V.hi.s5, inout);
BLAKE_G(x, 0x0E, V.lo.s3, V.lo.s4, V.hi.s1, V.hi.s6, inout);
}

V.lo ^= V.hi ^ tmpblock;

((uint8*)inout)[0]=V.lo;

((uint8*)inout)[0] = V.lo;
}


static __forceinline__ __host__ void Blake2Shost(uint32_t * inout, const uint32_t * inkey)
{
uint16 V;
Expand Down Expand Up @@ -489,9 +613,9 @@ __global__ __launch_bounds__(128, 1) void neoscrypt_gpu_hash_k01(int threads, ui
for (int i = 0; i < 128; ++i)
{
neoscrypt_chacha(X);
((ulonglong16 *)(W + shift))[i+1] = ((ulonglong16 *)X)[0];
// ((ulonglong16 *)(W + shift))[i+1] = ((ulonglong16 *)X)[0];

// ((uintx64 *)(W + shift))[i + 1] = ((uintx64 *)X)[0];
((uintx64 *)(W + shift))[i + 1] = ((uintx64 *)X)[0];
}


Expand Down Expand Up @@ -536,8 +660,8 @@ __global__ __launch_bounds__(128, 1) void neoscrypt_gpu_hash_k3(int threads, uin
for (int i = 0; i < 128; ++i)
{
neoscrypt_salsa(Z);
((ulonglong16 *)(W + shift))[i+1] = ((ulonglong16 *)Z)[0];
// ((uintx64 *)(W + shift))[i + 1] = ((uintx64 *)Z)[0];
// ((ulonglong16 *)(W + shift))[i+1] = ((ulonglong16 *)Z)[0];
((uintx64 *)(W + shift))[i + 1] = ((uintx64 *)Z)[0];
}


Expand Down

0 comments on commit 817bcae

Please sign in to comment.