Skip to content

Commit

Permalink
x11: echo and cubehash optimization
Browse files Browse the repository at this point in the history 
echo : 40.056ms -> 39.241ms
cube : 14.490ms -> 13.511ms

cube hash change look like useless (__device__ code in generally inlined)
but the reality proves that cuda documentation is wrong...

tpruvot: fixed dos lines ending in echo,
and used my style for cuda function attributes
  • Loading branch information
@sp-hash @tpruvot
sp-hash authored and tpruvot committed Nov 6, 2014
1 parent 12fafd5 commit 5be6811
Show file tree
Hide file tree
Showing 3 changed files with 71 additions and 104 deletions.
104 changes: 34 additions & 70 deletions x11/cuda_x11_aes.cu
View file
Expand Up @@ -319,49 +319,32 @@ static void aes_round(
uint32_t x0, uint32_t x1, uint32_t x2, uint32_t x3, uint32_t k0,
uint32_t &y0, uint32_t &y1, uint32_t &y2, uint32_t &y3)
{
uint32_t idx0, idx1, idx2, idx3;

idx0 = __byte_perm(x0, 0, 0x4440);
idx1 = __byte_perm(x1, 0, 0x4441) + 256;
idx2 = __byte_perm(x2, 0, 0x4442) + 512;
idx3 = __byte_perm(x3, 0, 0x4443) + 768;
y0 = xor4_32(
sharedMemory[idx0],
sharedMemory[idx1],
sharedMemory[idx2],
sharedMemory[idx3]);
y0 ^= k0;

idx0 = __byte_perm(x1, 0, 0x4440);
idx1 = __byte_perm(x2, 0, 0x4441) + 256;
idx2 = __byte_perm(x3, 0, 0x4442) + 512;
idx3 = __byte_perm(x0, 0, 0x4443) + 768;
sharedMemory[__byte_perm(x0, 0, 0x4440)],
sharedMemory[__byte_perm(x1, 0, 0x4441) + 256],
sharedMemory[__byte_perm(x2, 0, 0x4442) + 512],
sharedMemory[__byte_perm(x3, 0, 0x4443) + 768]);

y1 = xor4_32(
sharedMemory[idx0],
sharedMemory[idx1],
sharedMemory[idx2],
sharedMemory[idx3]);
sharedMemory[__byte_perm(x1, 0, 0x4440)],
sharedMemory[__byte_perm(x2, 0, 0x4441) + 256],
sharedMemory[__byte_perm(x3, 0, 0x4442) + 512],
sharedMemory[__byte_perm(x0, 0, 0x4443) + 768]);

idx0 = __byte_perm(x2, 0, 0x4440);
idx1 = __byte_perm(x3, 0, 0x4441) + 256;
idx2 = __byte_perm(x0, 0, 0x4442) + 512;
idx3 = __byte_perm(x1, 0, 0x4443) + 768;
y2 = xor4_32(
sharedMemory[idx0],
sharedMemory[idx1],
sharedMemory[idx2],
sharedMemory[idx3]); // ^k2
sharedMemory[__byte_perm(x2, 0, 0x4440)],
sharedMemory[__byte_perm(x3, 0, 0x4441) + 256],
sharedMemory[__byte_perm(x0, 0, 0x4442) + 512],
sharedMemory[__byte_perm(x1, 0, 0x4443) + 768]); // ^k2

y0 ^= k0;

idx0 = __byte_perm(x3, 0, 0x4440);
idx1 = __byte_perm(x0, 0, 0x4441) + 256;
idx2 = __byte_perm(x1, 0, 0x4442) + 512;
idx3 = __byte_perm(x2, 0, 0x4443) + 768;
y3 = xor4_32(
sharedMemory[idx0],
sharedMemory[idx1],
sharedMemory[idx2],
sharedMemory[idx3]); // ^k3
sharedMemory[__byte_perm(x3, 0, 0x4440)],
sharedMemory[__byte_perm(x0, 0, 0x4441) + 256],
sharedMemory[__byte_perm(x1, 0, 0x4442) + 512],
sharedMemory[__byte_perm(x2, 0, 0x4443) + 768]); // ^k3
}

__device__
Expand All @@ -370,46 +353,27 @@ static void aes_round(
uint32_t x0, uint32_t x1, uint32_t x2, uint32_t x3,
uint32_t &y0, uint32_t &y1, uint32_t &y2, uint32_t &y3)
{
uint32_t idx0, idx1, idx2, idx3;

idx0 = __byte_perm(x0, 0, 0x4440);
idx1 = __byte_perm(x1, 0, 0x4441) + 256;
idx2 = __byte_perm(x2, 0, 0x4442) + 512;
idx3 = __byte_perm(x3, 0, 0x4443) + 768;
y0 = xor4_32(
sharedMemory[idx0],
sharedMemory[idx1],
sharedMemory[idx2],
sharedMemory[idx3]);

idx0 = __byte_perm(x1, 0, 0x4440);
idx1 = __byte_perm(x2, 0, 0x4441) + 256;
idx2 = __byte_perm(x3, 0, 0x4442) + 512;
idx3 = __byte_perm(x0, 0, 0x4443) + 768;
sharedMemory[__byte_perm(x0, 0, 0x4440)],
sharedMemory[__byte_perm(x1, 0, 0x4441) + 256],
sharedMemory[__byte_perm(x2, 0, 0x4442) + 512],
sharedMemory[__byte_perm(x3, 0, 0x4443) + 768]);

y1 = xor4_32(
sharedMemory[idx0],
sharedMemory[idx1],
sharedMemory[idx2],
sharedMemory[idx3]);
sharedMemory[__byte_perm(x1, 0, 0x4440)],
sharedMemory[__byte_perm(x2, 0, 0x4441) + 256],
sharedMemory[__byte_perm(x3, 0, 0x4442) + 512],
sharedMemory[__byte_perm(x0, 0, 0x4443) + 768]);

idx0 = __byte_perm(x2, 0, 0x4440);
idx1 = __byte_perm(x3, 0, 0x4441) + 256;
idx2 = __byte_perm(x0, 0, 0x4442) + 512;
idx3 = __byte_perm(x1, 0, 0x4443) + 768;
y2 = xor4_32(
sharedMemory[idx0],
sharedMemory[idx1],
sharedMemory[idx2],
sharedMemory[idx3]); // ^k2
sharedMemory[__byte_perm(x2, 0, 0x4440)],
sharedMemory[__byte_perm(x3, 0, 0x4441) + 256],
sharedMemory[__byte_perm(x0, 0, 0x4442) + 512],
sharedMemory[__byte_perm(x1, 0, 0x4443) + 768]); // ^k2

idx0 = __byte_perm(x3, 0, 0x4440);
idx1 = __byte_perm(x0, 0, 0x4441) + 256;
idx2 = __byte_perm(x1, 0, 0x4442) + 512;
idx3 = __byte_perm(x2, 0, 0x4443) + 768;
y3 = xor4_32(
sharedMemory[idx0],
sharedMemory[idx1],
sharedMemory[idx2],
sharedMemory[idx3]); // ^k3
sharedMemory[__byte_perm(x3, 0, 0x4440)],
sharedMemory[__byte_perm(x0, 0, 0x4441) + 256],
sharedMemory[__byte_perm(x1, 0, 0x4442) + 512],
sharedMemory[__byte_perm(x2, 0, 0x4443) + 768]); // ^k3
}
24 changes: 15 additions & 9 deletions x11/cuda_x11_cubehash512.cu
View file
Expand Up @@ -34,7 +34,8 @@ static const uint32_t c_IV_512[32] = {
0x7795D246, 0xD43E3B44
};

static __device__ void rrounds(uint32_t x[2][2][2][2][2])
__device__ __forceinline__
static void rrounds(uint32_t x[2][2][2][2][2])
{
int r;
int j;
Expand Down Expand Up @@ -150,8 +151,8 @@ static __device__ void rrounds(uint32_t x[2][2][2][2][2])
}
}


static __device__ void block_tox(uint32_t block[16], uint32_t x[2][2][2][2][2])
__device__ __forceinline__
static void block_tox(uint32_t block[16], uint32_t x[2][2][2][2][2])
{
int k;
int l;
Expand All @@ -167,7 +168,8 @@ static __device__ void block_tox(uint32_t block[16], uint32_t x[2][2][2][2][2])
x[0][0][k][l][m] ^= *in++;
}

static __device__ void hash_fromx(uint32_t hash[16], uint32_t x[2][2][2][2][2])
__device__ __forceinline__
static void hash_fromx(uint32_t hash[16], uint32_t x[2][2][2][2][2])
{
int j;
int k;
Expand All @@ -186,7 +188,8 @@ static __device__ void hash_fromx(uint32_t hash[16], uint32_t x[2][2][2][2][2])
*out++ = x[0][j][k][l][m];
}

void __device__ Init(uint32_t x[2][2][2][2][2])
__device__
void Init(uint32_t x[2][2][2][2][2])
{
int i,j,k,l,m;
#if 0
Expand Down Expand Up @@ -227,15 +230,17 @@ void __device__ Init(uint32_t x[2][2][2][2][2])
#endif
}

void __device__ Update32(uint32_t x[2][2][2][2][2], const BitSequence *data)
__device__ __forceinline__
void Update32(uint32_t x[2][2][2][2][2], const BitSequence *data)
{
/* "xor the block into the first b bytes of the state" */
/* "and then transform the state invertibly through r identical rounds" */
block_tox((uint32_t*)data, x);
rrounds(x);
}

void __device__ Final(uint32_t x[2][2][2][2][2], BitSequence *hashval)
__device__ __forceinline__
void Final(uint32_t x[2][2][2][2][2], BitSequence *hashval)
{
int i;

Expand All @@ -252,8 +257,9 @@ void __device__ Final(uint32_t x[2][2][2][2][2], BitSequence *hashval)


/***************************************************/
// Die Hash-Funktion
__global__ void x11_cubehash512_gpu_hash_64(int threads, uint32_t startNounce, uint64_t *g_hash, uint32_t *g_nonceVector)
// GPU Hash Function
__global__
void x11_cubehash512_gpu_hash_64(int threads, uint32_t startNounce, uint64_t *g_hash, uint32_t *g_nonceVector)
{
int thread = (blockDim.x * blockIdx.x + threadIdx.x);
if (thread < threads)
Expand Down
47 changes: 22 additions & 25 deletions x11/cuda_x11_echo.cu
View file
Expand Up @@ -75,41 +75,38 @@ __device__ __forceinline__ void cuda_echo_round(
}

// Mix Columns
#pragma unroll 4
for(int i=0;i<4;i++) // Schleife über je 2*uint32_t
#pragma unroll
for (int i = 0; i<4; i++) // Schleife über je 2*uint32_t
{
#pragma unroll 4
for(int j=0;j<4;j++) // Schleife über die elemnte
#pragma unroll 64
for (int idx = 0; idx<64; idx += 16) // Schleife über die elemnte
{
int idx = j<<2; // j*4

uint32_t a = W[ ((idx + 0)<<2) + i];
uint32_t b = W[ ((idx + 1)<<2) + i];
uint32_t c = W[ ((idx + 2)<<2) + i];
uint32_t d = W[ ((idx + 3)<<2) + i];
uint32_t a = W[idx + i];
uint32_t b = W[idx + i + 4];
uint32_t c = W[idx + i + 8];
uint32_t d = W[idx + i + 12];

uint32_t ab = a ^ b;
uint32_t bc = b ^ c;
uint32_t cd = c ^ d;

uint32_t t;
t = ((ab & 0x80808080) >> 7);
uint32_t abx = t<<4 ^ t<<3 ^ t<<1 ^ t;
t = ((bc & 0x80808080) >> 7);
uint32_t bcx = t<<4 ^ t<<3 ^ t<<1 ^ t;
t = ((cd & 0x80808080) >> 7);
uint32_t cdx = t<<4 ^ t<<3 ^ t<<1 ^ t;

abx ^= ((ab & 0x7F7F7F7F) << 1);
bcx ^= ((bc & 0x7F7F7F7F) << 1);
cdx ^= ((cd & 0x7F7F7F7F) << 1);

W[ ((idx + 0)<<2) + i] = abx ^ bc ^ d;
W[ ((idx + 1)<<2) + i] = bcx ^ a ^ cd;
W[ ((idx + 2)<<2) + i] = cdx ^ ab ^ d;
W[ ((idx + 3)<<2) + i] = abx ^ bcx ^ cdx ^ ab ^ c;
uint32_t t, t2, t3;
t = (ab & 0x80808080);
t2 = (bc & 0x80808080);
t3 = (cd & 0x80808080);

uint32_t abx = (t >> 7) * 27 ^ ((ab^t) << 1);
uint32_t bcx = (t2 >> 7) * 27 ^ ((bc^t2) << 1);
uint32_t cdx = (t3 >> 7) * 27 ^ ((cd^t3) << 1);

W[idx + i] = abx ^ bc ^ d;
W[idx + i + 4] = bcx ^ a ^ cd;
W[idx + i + 8] = cdx ^ ab ^ d;
W[idx + i + 12] = abx ^ bcx ^ cdx ^ ab ^ c;
}
}

}

__global__ void x11_echo512_gpu_hash_64(int threads, uint32_t startNounce, uint64_t *g_hash, uint32_t *g_nonceVector)
Expand Down

0 comments on commit 5be6811

Please sign in to comment.