Browse Source

jh512: import and improve klaus and sp changes

did not import the extra final function, which should stay compatible
with the common cuda_check_hash()
master
Tanguy Pruvot 10 years ago
parent
commit
9734186a37
  1. 294
      quark/cuda_jh512.cu

294
quark/cuda_jh512.cu

@ -1,25 +1,6 @@ @@ -1,25 +1,6 @@
#include "cuda_helper.h"
typedef struct {
uint32_t x[8][4]; /*the 1024-bit state, ( x[i][0] || x[i][1] || x[i][2] || x[i][3] ) is the ith row of the state in the pseudocode*/
uint32_t buffer[16]; /*the 512-bit message block to be hashed;*/
} hashState;
/*42 round constants, each round constant is 32-byte (256-bit)*/
__constant__ uint32_t c_INIT_bitslice[8][4];
__constant__ unsigned char c_E8_bitslice_roundconstant[42][32];
const uint32_t h_INIT_bitslice[8][4] = {
{ 0x964bd16f, 0x17aa003e, 0x052e6a63, 0x43d5157a},
{ 0x8d5e228a, 0x0bef970c, 0x591234e9, 0x61c3b3f2},
{ 0xc1a01d89, 0x1e806f53, 0x6b05a92a, 0x806d2bea},
{ 0xdbcc8e58, 0xa6ba7520, 0x763a0fa9, 0xf73bf8ba},
{ 0x05e66901, 0x694ae341, 0x8e8ab546, 0x5ae66f2e},
{ 0xd0a74710, 0x243c84c1, 0xb1716e3b, 0x99c15a2d},
{ 0xecf657cf, 0x56f8b19d, 0x7c8806a7, 0x56b11657},
{ 0xdffcc2e3, 0xfb1785e6, 0x78465a54, 0x4bdd8ccc} };
const unsigned char h_E8_bitslice_roundconstant[42][32]={
__constant__ unsigned char c_E8_bitslice_roundconstant[42][32] = {
{ 0x72, 0xd5, 0xde, 0xa2, 0xdf, 0x15, 0xf8, 0x67, 0x7b, 0x84, 0x15, 0xa, 0xb7, 0x23, 0x15, 0x57, 0x81, 0xab, 0xd6, 0x90, 0x4d, 0x5a, 0x87, 0xf6, 0x4e, 0x9f, 0x4f, 0xc5, 0xc3, 0xd1, 0x2b, 0x40 },
{ 0xea, 0x98, 0x3a, 0xe0, 0x5c, 0x45, 0xfa, 0x9c, 0x3, 0xc5, 0xd2, 0x99, 0x66, 0xb2, 0x99, 0x9a, 0x66, 0x2, 0x96, 0xb4, 0xf2, 0xbb, 0x53, 0x8a, 0xb5, 0x56, 0x14, 0x1a, 0x88, 0xdb, 0xa2, 0x31 },
{ 0x3, 0xa3, 0x5a, 0x5c, 0x9a, 0x19, 0xe, 0xdb, 0x40, 0x3f, 0xb2, 0xa, 0x87, 0xc1, 0x44, 0x10, 0x1c, 0x5, 0x19, 0x80, 0x84, 0x9e, 0x95, 0x1d, 0x6f, 0x33, 0xeb, 0xad, 0x5e, 0xe7, 0xcd, 0xdc },
@ -65,25 +46,24 @@ const unsigned char h_E8_bitslice_roundconstant[42][32]={ @@ -65,25 +46,24 @@ const unsigned char h_E8_bitslice_roundconstant[42][32]={
#define SWAP4(x,y)\
y = (x & 0xf0f0f0f0UL); \
x ^= y; \
y >>= 4; \
x <<= 4; \
x |= y;
x = (x ^ y); \
y = (y >> 4); \
x = (x << 4); \
x= x | y;
#define SWAP2(x,y)\
y = (x & 0xccccccccUL); \
x ^= y; \
y >>= 2; \
x <<= 2; \
x |= y;
x = (x ^ y); \
y = (y >> 2); \
x = (x << 2); \
x= x | y;
#define SWAP1(x,y)\
y = (x & 0xaaaaaaaaUL); \
x ^= y; \
y >>= 1; \
x += x; \
x |= y;
x = (x ^ y); \
y = (y >> 1); \
x = x + x; \
x= x | y;
/*swapping bits 16i||16i+1||......||16i+7 with bits 16i+8||16i+9||......||16i+15 of 32-bit x*/
//#define SWAP8(x) (x) = ((((x) & 0x00ff00ffUL) << 8) | (((x) & 0xff00ff00UL) >> 8));
#define SWAP8(x) (x) = __byte_perm(x, x, 0x2301);
@ -116,7 +96,7 @@ const unsigned char h_E8_bitslice_roundconstant[42][32]={ @@ -116,7 +96,7 @@ const unsigned char h_E8_bitslice_roundconstant[42][32]={
m1 ^= (temp0 & (m0)); \
m2 ^= temp0;
__device__ __forceinline__ void Sbox_and_MDS_layer(hashState* state, uint32_t roundnumber)
static __device__ __forceinline__ void Sbox_and_MDS_layer(uint32_t x[8][4], uint32_t roundnumber)
{
uint32_t temp0;
uint32_t cc0, cc1;
@ -125,254 +105,198 @@ __device__ __forceinline__ void Sbox_and_MDS_layer(hashState* state, uint32_t ro @@ -125,254 +105,198 @@ __device__ __forceinline__ void Sbox_and_MDS_layer(hashState* state, uint32_t ro
for (int i = 0; i < 4; i++) {
cc0 = ((uint32_t*)c_E8_bitslice_roundconstant[roundnumber])[i];
cc1 = ((uint32_t*)c_E8_bitslice_roundconstant[roundnumber])[i + 4];
Sbox(state->x[0][i],state->x[2][i], state->x[4][i], state->x[6][i], cc0);
Sbox(state->x[1][i],state->x[3][i], state->x[5][i], state->x[7][i], cc1);
L(state->x[0][i],state->x[2][i],state->x[4][i],state->x[6][i],state->x[1][i],state->x[3][i],state->x[5][i],state->x[7][i]);
Sbox(x[0][i], x[2][i], x[4][i], x[6][i], cc0);
Sbox(x[1][i], x[3][i], x[5][i], x[7][i], cc1);
L(x[0][i], x[2][i], x[4][i], x[6][i], x[1][i], x[3][i], x[5][i], x[7][i]);
}
}
__device__ __forceinline__ void RoundFunction0(hashState* state, uint32_t roundnumber)
static __device__ __forceinline__ void RoundFunction0(uint32_t x[8][4], uint32_t roundnumber)
{
Sbox_and_MDS_layer(state, roundnumber);
Sbox_and_MDS_layer(x, roundnumber);
#pragma unroll 4
for (int j = 1; j < 8; j = j + 2)
{
uint32_t y;
SWAP1(state->x[j][0], y);
SWAP1(state->x[j][1], y);
SWAP1(state->x[j][2], y);
SWAP1(state->x[j][3], y);
SWAP1(x[j][0], y);
SWAP1(x[j][1], y);
SWAP1(x[j][2], y);
SWAP1(x[j][3], y);
}
}
__device__ __forceinline__ void RoundFunction1(hashState* state, uint32_t roundnumber)
static __device__ __forceinline__ void RoundFunction1(uint32_t x[8][4], uint32_t roundnumber)
{
Sbox_and_MDS_layer(state, roundnumber);
Sbox_and_MDS_layer(x, roundnumber);
#pragma unroll 4
for (int j = 1; j < 8; j = j + 2)
{
uint32_t y;
SWAP2(state->x[j][0], y);
SWAP2(state->x[j][1], y);
SWAP2(state->x[j][2], y);
SWAP2(state->x[j][3], y);
SWAP2(x[j][0], y);
SWAP2(x[j][1], y);
SWAP2(x[j][2], y);
SWAP2(x[j][3], y);
}
}
__device__ __forceinline__ void RoundFunction2(hashState* state, uint32_t roundnumber)
static __device__ __forceinline__ void RoundFunction2(uint32_t x[8][4], uint32_t roundnumber)
{
Sbox_and_MDS_layer(state, roundnumber);
Sbox_and_MDS_layer(x, roundnumber);
#pragma unroll 4
for (int j = 1; j < 8; j = j + 2)
{
uint32_t y;
SWAP4(state->x[j][0], y);
SWAP4(state->x[j][1], y);
SWAP4(state->x[j][2], y);
SWAP4(state->x[j][3], y);
SWAP4(x[j][0], y);
SWAP4(x[j][1], y);
SWAP4(x[j][2], y);
SWAP4(x[j][3], y);
}
}
__device__ __forceinline__ void RoundFunction3(hashState* state, uint32_t roundnumber)
static __device__ __forceinline__ void RoundFunction3(uint32_t x[8][4], uint32_t roundnumber)
{
Sbox_and_MDS_layer(state, roundnumber);
Sbox_and_MDS_layer(x, roundnumber);
#pragma unroll 4
for (int j = 1; j < 8; j = j + 2)
{
#pragma unroll 4
for (int i = 0; i < 4; i++) SWAP8(state->x[j][i]);
for (int i = 0; i < 4; i++) SWAP8(x[j][i]);
}
}
__device__ __forceinline__ void RoundFunction4(hashState* state, uint32_t roundnumber)
static __device__ __forceinline__ void RoundFunction4(uint32_t x[8][4], uint32_t roundnumber)
{
Sbox_and_MDS_layer(state, roundnumber);
Sbox_and_MDS_layer(x, roundnumber);
#pragma unroll 4
for (int j = 1; j < 8; j = j + 2)
{
#pragma unroll 4
for (int i = 0; i < 4; i++) SWAP16(state->x[j][i]);
for (int i = 0; i < 4; i++) SWAP16(x[j][i]);
}
}
__device__ __forceinline__ void RoundFunction5(hashState* state, uint32_t roundnumber)
static __device__ __forceinline__ void RoundFunction5(uint32_t x[8][4], uint32_t roundnumber)
{
uint32_t temp0;
Sbox_and_MDS_layer(state, roundnumber);
Sbox_and_MDS_layer(x, roundnumber);
#pragma unroll 4
for (int j = 1; j < 8; j = j + 2)
{
#pragma unroll 2
for (int i = 0; i < 4; i = i + 2) {
temp0 = state->x[j][i]; state->x[j][i] = state->x[j][i+1]; state->x[j][i+1] = temp0;
temp0 = x[j][i]; x[j][i] = x[j][i + 1]; x[j][i + 1] = temp0;
}
}
}
__device__ __forceinline__ void RoundFunction6(hashState* state, uint32_t roundnumber)
static __device__ __forceinline__ void RoundFunction6(uint32_t x[8][4], uint32_t roundnumber)
{
uint32_t temp0;
Sbox_and_MDS_layer(state, roundnumber);
Sbox_and_MDS_layer(x, roundnumber);
#pragma unroll 4
for (int j = 1; j < 8; j = j + 2)
{
#pragma unroll 2
for (int i = 0; i < 2; i++) {
temp0 = state->x[j][i]; state->x[j][i] = state->x[j][i+2]; state->x[j][i+2] = temp0;
temp0 = x[j][i]; x[j][i] = x[j][i + 2]; x[j][i + 2] = temp0;
}
}
}
/*The bijective function E8, in bitslice form */
__device__ __forceinline__ void E8(hashState *state)
static __device__ __forceinline__ void E8(uint32_t x[8][4])
{
/*perform 6 rounds*/
//#pragma unroll 6
for (int i = 0; i < 42; i += 7)
{
RoundFunction0(state, i);
RoundFunction1(state, i+1);
RoundFunction2(state, i+2);
RoundFunction3(state, i+3);
RoundFunction4(state, i+4);
RoundFunction5(state, i+5);
RoundFunction6(state, i+6);
RoundFunction0(x, i);
RoundFunction1(x, i + 1);
RoundFunction2(x, i + 2);
RoundFunction3(x, i + 3);
RoundFunction4(x, i + 4);
RoundFunction5(x, i + 5);
RoundFunction6(x, i + 6);
}
}
/*The compression function F8 */
__device__ __forceinline__ void F8(hashState *state)
{
/*xor the 512-bit message with the fist half of the 1024-bit hash state*/
#pragma unroll 16
for (int i = 0; i < 16; i++) state->x[i >> 2][i & 3] ^= ((uint32_t*)state->buffer)[i];
/*the bijective function E8 */
E8(state);
/*xor the 512-bit message with the second half of the 1024-bit hash state*/
#pragma unroll 16
for (int i = 0; i < 16; i++) state->x[(16+i) >> 2][(16+i) & 3] ^= ((uint32_t*)state->buffer)[i];
}
__device__ __forceinline__ void JHHash(const uint32_t *data, uint32_t *hashval)
// Die Hash-Funktion
__global__ __launch_bounds__(256, 4)
void quark_jh512_gpu_hash_64(uint32_t threads, uint32_t startNounce, uint32_t *const __restrict__ g_hash, const uint32_t *const __restrict__ g_nonceVector)
{
hashState state;
/*load the intital hash value H0 into state*/
/*
#define INIT(a,b,c,d) ((a) | ((b)<<8) | ((c)<<16) | ((d)<<24))
state.x[0][0] = INIT(0x6f,0xd1,0x4b,0x96);
state.x[0][1] = INIT(0x3e,0x00,0xaa,0x17);
state.x[0][2] = INIT(0x63,0x6a,0x2e,0x05);
state.x[0][3] = INIT(0x7a,0x15,0xd5,0x43);
state.x[1][0] = INIT(0x8a,0x22,0x5e,0x8d);
state.x[1][1] = INIT(0x0c,0x97,0xef,0x0b);
state.x[1][2] = INIT(0xe9,0x34,0x12,0x59);
state.x[1][3] = INIT(0xf2,0xb3,0xc3,0x61);
state.x[2][0] = INIT(0x89,0x1d,0xa0,0xc1);
state.x[2][1] = INIT(0x53,0x6f,0x80,0x1e);
state.x[2][2] = INIT(0x2a,0xa9,0x05,0x6b);
state.x[2][3] = INIT(0xea,0x2b,0x6d,0x80);
state.x[3][0] = INIT(0x58,0x8e,0xcc,0xdb);
state.x[3][1] = INIT(0x20,0x75,0xba,0xa6);
state.x[3][2] = INIT(0xa9,0x0f,0x3a,0x76);
state.x[3][3] = INIT(0xba,0xf8,0x3b,0xf7);
state.x[4][0] = INIT(0x01,0x69,0xe6,0x05);
state.x[4][1] = INIT(0x41,0xe3,0x4a,0x69);
state.x[4][2] = INIT(0x46,0xb5,0x8a,0x8e);
state.x[4][3] = INIT(0x2e,0x6f,0xe6,0x5a);
state.x[5][0] = INIT(0x10,0x47,0xa7,0xd0);
state.x[5][1] = INIT(0xc1,0x84,0x3c,0x24);
state.x[5][2] = INIT(0x3b,0x6e,0x71,0xb1);
state.x[5][3] = INIT(0x2d,0x5a,0xc1,0x99);
state.x[6][0] = INIT(0xcf,0x57,0xf6,0xec);
state.x[6][1] = INIT(0x9d,0xb1,0xf8,0x56);
state.x[6][2] = INIT(0xa7,0x06,0x88,0x7c);
state.x[6][3] = INIT(0x57,0x16,0xb1,0x56);
state.x[7][0] = INIT(0xe3,0xc2,0xfc,0xdf);
state.x[7][1] = INIT(0xe6,0x85,0x17,0xfb);
state.x[7][2] = INIT(0x54,0x5a,0x46,0x78);
state.x[7][3] = INIT(0xcc,0x8c,0xdd,0x4b);
*/
#pragma unroll 8
for(int j=0;j<8;j++)
const uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x);
if (thread < threads)
{
#pragma unroll 4
for(int i=0;i<4;i++)
state.x[j][i] = c_INIT_bitslice[j][i];
}
const uint32_t nounce = (g_nonceVector != NULL) ? g_nonceVector[thread] : (startNounce + thread);
const uint32_t hashPosition = nounce - startNounce;
uint32_t *const Hash = &g_hash[hashPosition * 16U];
uint32_t x[8][4] = {
{ 0x964bd16f, 0x17aa003e, 0x052e6a63, 0x43d5157a },
{ 0x8d5e228a, 0x0bef970c, 0x591234e9, 0x61c3b3f2 },
{ 0xc1a01d89, 0x1e806f53, 0x6b05a92a, 0x806d2bea },
{ 0xdbcc8e58, 0xa6ba7520, 0x763a0fa9, 0xf73bf8ba },
{ 0x05e66901, 0x694ae341, 0x8e8ab546, 0x5ae66f2e },
{ 0xd0a74710, 0x243c84c1, 0xb1716e3b, 0x99c15a2d },
{ 0xecf657cf, 0x56f8b19d, 0x7c8806a7, 0x56b11657 },
{ 0xdffcc2e3, 0xfb1785e6, 0x78465a54, 0x4bdd8ccc }
};
#pragma unroll 16
for (int i=0; i < 16; ++i) state.buffer[i] = data[i];
F8(&state);
for (int i = 0; i < 16; i++)
x[i >> 2][i & 3] ^= Hash[i];
/*pad the message when databitlen is multiple of 512 bits, then process the padded block*/
state.buffer[0] = 0x80;
#pragma unroll 14
for (int i=1; i < 15; i++) state.buffer[i] = 0;
state.buffer[15] = 0x00020000;
F8(&state);
E8(x);
/*truncating the final hash value to generate the message digest*/
#pragma unroll 16
for (int i=0; i < 16; ++i) hashval[i] = state.x[4][i];
}
// Die Hash-Funktion
__global__ __launch_bounds__(256, 3)
void quark_jh512_gpu_hash_64(uint32_t threads, uint32_t startNounce, uint64_t *g_hash, uint32_t *g_nonceVector)
{
uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x);
if (thread < threads)
{
uint32_t nounce = (g_nonceVector != NULL) ? g_nonceVector[thread] : (startNounce + thread);
int hashPosition = nounce - startNounce;
uint32_t *Hash = (uint32_t*)&g_hash[8 * hashPosition];
JHHash(Hash, Hash);
for (int i = 0; i < 16; i++)
x[(16 + i) >> 2][(16 + i) & 3] ^= Hash[i];
x[0 >> 2][0 & 3] ^= 0x80;
x[15 >> 2][15 & 3] ^= 0x00020000;
E8(x);
x[(16 + 0) >> 2][(16 + 0) & 3] ^= 0x80;
x[(16 + 15) >> 2][(16 + 15) & 3] ^= 0x00020000;
Hash[0] = x[4][0];
Hash[1] = x[4][1];
Hash[2] = x[4][2];
Hash[3] = x[4][3];
Hash[4] = x[5][0];
Hash[5] = x[5][1];
Hash[6] = x[5][2];
Hash[7] = x[5][3];
Hash[8] = x[6][0];
Hash[9] = x[6][1];
Hash[10] = x[6][2];
Hash[11] = x[6][3];
Hash[12] = x[7][0];
Hash[13] = x[7][1];
Hash[14] = x[7][2];
Hash[15] = x[7][3];
}
}
// Setup-Funktionen
__host__ void quark_jh512_cpu_init(int thr_id, uint32_t threads)
{
cudaMemcpyToSymbol( c_E8_bitslice_roundconstant,
h_E8_bitslice_roundconstant,
sizeof(h_E8_bitslice_roundconstant),
0, cudaMemcpyHostToDevice);
cudaMemcpyToSymbol( c_INIT_bitslice,
h_INIT_bitslice,
sizeof(h_INIT_bitslice),
0, cudaMemcpyHostToDevice);
}
__host__ void quark_jh512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order)
__host__
void quark_jh512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order)
{
const uint32_t threadsperblock = 256;
// berechne wie viele Thread Blocks wir brauchen
dim3 grid((threads + threadsperblock-1)/threadsperblock);
dim3 block(threadsperblock);
// Größe des dynamischen Shared Memory Bereichs
size_t shared_size = 0;
quark_jh512_gpu_hash_64<<<grid, block, shared_size>>>(threads, startNounce, (uint64_t*)d_hash, d_nonceVector);
MyStreamSynchronize(NULL, order, thr_id);
quark_jh512_gpu_hash_64<<<grid, block>>>(threads, startNounce, d_hash, d_nonceVector);
}
// Setup function
__host__ void quark_jh512_cpu_init(int thr_id, uint32_t threads) {}

Loading…
Cancel
Save