mirror of
https://github.com/hashcat/hashcat.git
synced 2025-02-16 17:42:04 +00:00
Improved SAP CODVN F/G password cracking performance by working around byte endian swaps
This commit is contained in:
jsteube
parent
899413f88b
commit
425031c9c3
@ -17,46 +17,42 @@
|
||||
#include "inc_rp.cl"
|
||||
#include "inc_simd.cl"
|
||||
|
||||
#define GETSHIFTEDINT(a,n) amd_bytealign ((a)[((n)/4)+1], (a)[((n)/4)+0], (n))
|
||||
|
||||
#define SETSHIFTEDINT(a,n,v) \
|
||||
{ \
|
||||
const u32 s = ((n) & 3) * 8; \
|
||||
const u64 x = (u64) (v) << s; \
|
||||
(a)[((n)/4)+0] |= x; \
|
||||
(a)[((n)/4)+1] = x >> 32; \
|
||||
}
|
||||
|
||||
__constant u32 theMagicArray[64] =
|
||||
{
|
||||
0x1451ac91,0x4354679f,0xe03be724,0xc27b7428,0xeb133386,0x5ccb4f5a,0x37730a08,0x2f1c5d0e,
|
||||
0xe5e68f33,0xddae9bf8,0x8d4bf216,0xdcd4e12c,0x9ddfcbb0,0x176d70d4,0x3f424df9,0x94111b9b,
|
||||
0x9bc15b9f,0x039d0506,0x8a135e9d,0xe86a9a1e,0x17147cd9,0xf62ac758,0x0a6399a1,0xc370fdd7,
|
||||
0x13745ef6,0x040bc903,0x26f79826,0x2593928a,0x230da2b0,0x6d7963ed,0x3cfa3213,0xa39a0235,
|
||||
0x0a8eddb3,0xc351bf24,0x9f55cd7c,0x4c94af37,0x82520829,0x374e3bb2,0x9107179f,0xcdfd3b11,
|
||||
0, 0, 0, 0, 0, 0, 0, 0,
|
||||
0, 0, 0, 0, 0, 0, 0, 0,
|
||||
0, 0, 0, 0, 0, 0, 0, 0
|
||||
0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f,
|
||||
0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194,
|
||||
0x9f5bc19b, 0x06059d03, 0x9d5e138a, 0x1e9a6ae8, 0xd97c1417, 0x58c72af6, 0xa199630a, 0xd7fd70c3,
|
||||
0xf65e7413, 0x03c90b04, 0x2698f726, 0x8a929325, 0xb0a20d23, 0xed63796d, 0x1332fa3c, 0x35029aa3,
|
||||
0xb3dd8e0a, 0x24bf51c3, 0x7ccd559f, 0x37af944c, 0x29085282, 0xb23b4e37, 0x9f170791, 0x113bfdcd,
|
||||
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
|
||||
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
|
||||
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
|
||||
};
|
||||
|
||||
void swap_buffer (u32 final[16])
|
||||
static u32 GETSHIFTEDINT_CONST (__constant u32 *a, const int n)
|
||||
{
|
||||
final[ 0] = swap32 (final[ 0]);
|
||||
final[ 1] = swap32 (final[ 1]);
|
||||
final[ 2] = swap32 (final[ 2]);
|
||||
final[ 3] = swap32 (final[ 3]);
|
||||
final[ 4] = swap32 (final[ 4]);
|
||||
final[ 5] = swap32 (final[ 5]);
|
||||
final[ 6] = swap32 (final[ 6]);
|
||||
final[ 7] = swap32 (final[ 7]);
|
||||
final[ 8] = swap32 (final[ 8]);
|
||||
final[ 9] = swap32 (final[ 9]);
|
||||
final[10] = swap32 (final[10]);
|
||||
final[11] = swap32 (final[11]);
|
||||
final[12] = swap32 (final[12]);
|
||||
final[13] = swap32 (final[13]);
|
||||
final[14] = swap32 (final[14]);
|
||||
final[15] = swap32 (final[15]);
|
||||
const int d = n / 4;
|
||||
const int m = n & 3;
|
||||
|
||||
u64 tmp = hl32_to_64_S (a[d + 0], a[d + 1]);
|
||||
|
||||
tmp <<= m * 8;
|
||||
|
||||
return h32_from_64_S (tmp);
|
||||
}
|
||||
|
||||
|
||||
static void SETSHIFTEDINT (u32 *a, const int n, const u32 v)
|
||||
{
|
||||
const int d = n / 4;
|
||||
const int m = n & 3;
|
||||
|
||||
u64 tmp = hl32_to_64_S (v, 0);
|
||||
|
||||
tmp >>= m * 8;
|
||||
|
||||
a[d + 0] |= h32_from_64_S (tmp);
|
||||
a[d + 1] = l32_from_64_S (tmp);
|
||||
}
|
||||
|
||||
void sha1_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[5])
|
||||
@ -299,6 +295,22 @@ __kernel void m07800_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
final[13] = swap32 (w3[1] | s3[1]);
|
||||
final[14] = 0;
|
||||
final[15] = pw_salt_len * 8;
|
||||
final[16] = 0;
|
||||
final[17] = 0;
|
||||
final[18] = 0;
|
||||
final[19] = 0;
|
||||
final[20] = 0;
|
||||
final[21] = 0;
|
||||
final[22] = 0;
|
||||
final[23] = 0;
|
||||
final[24] = 0;
|
||||
final[25] = 0;
|
||||
final[26] = 0;
|
||||
final[27] = 0;
|
||||
final[28] = 0;
|
||||
final[29] = 0;
|
||||
final[30] = 0;
|
||||
final[31] = 0;
|
||||
|
||||
u32 digest[5];
|
||||
|
||||
@ -344,21 +356,24 @@ __kernel void m07800_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
digest[3] = SHA1M_D;
|
||||
digest[4] = SHA1M_E;
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (int i = 0; i < 32; i++) final[i] = 0;
|
||||
final[ 0] = swap32_S (w0[0]);
|
||||
final[ 1] = swap32_S (w0[1]);
|
||||
final[ 2] = swap32_S (w0[2]);
|
||||
final[ 3] = swap32_S (w0[3]);
|
||||
final[ 4] = swap32_S (w1[0]);
|
||||
final[ 5] = swap32_S (w1[1]);
|
||||
final[ 6] = swap32_S (w1[2]);
|
||||
final[ 7] = swap32_S (w1[3]);
|
||||
final[ 8] = 0;
|
||||
final[ 9] = 0;
|
||||
final[10] = 0;
|
||||
final[11] = 0;
|
||||
final[12] = 0;
|
||||
final[13] = 0;
|
||||
final[14] = 0;
|
||||
final[15] = 0;
|
||||
|
||||
final[0] = w0[0];
|
||||
final[1] = w0[1];
|
||||
final[2] = w0[2];
|
||||
final[3] = w0[3];
|
||||
final[4] = w1[0];
|
||||
final[5] = w1[1];
|
||||
final[6] = w1[2];
|
||||
final[7] = w1[3];
|
||||
|
||||
u32 final_len = out_len;
|
||||
u32 final_len = pw_len;
|
||||
|
||||
int i;
|
||||
|
||||
@ -366,14 +381,14 @@ __kernel void m07800_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (i = 0; i < lengthMagicArray - 4; i += 4)
|
||||
{
|
||||
const u32 tmp = GETSHIFTEDINT (theMagicArray, offsetMagicArray + i);
|
||||
const u32 tmp = GETSHIFTEDINT_CONST (theMagicArray, offsetMagicArray + i);
|
||||
|
||||
SETSHIFTEDINT (final, final_len + i, tmp);
|
||||
}
|
||||
|
||||
const u32 mask = 0xffffffff >> (((i - lengthMagicArray) & 3) * 8);
|
||||
const u32 mask = 0xffffffff << (((4 - (lengthMagicArray - i)) & 3) * 8);
|
||||
|
||||
const u32 tmp = GETSHIFTEDINT (theMagicArray, offsetMagicArray + i) & mask;
|
||||
const u32 tmp = GETSHIFTEDINT_CONST (theMagicArray, offsetMagicArray + i) & mask;
|
||||
|
||||
SETSHIFTEDINT (final, final_len + i, tmp);
|
||||
|
||||
@ -383,7 +398,7 @@ __kernel void m07800_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (i = 0; i < salt_len + 1; i += 4) // +1 for the 0x80
|
||||
{
|
||||
const u32 tmp = salt_buf[i / 4]; // attention, int[] not char[]
|
||||
const u32 tmp = swap32_S (salt_buf[i / 4]); // attention, int[] not char[]
|
||||
|
||||
SETSHIFTEDINT (final, final_len + i, tmp);
|
||||
}
|
||||
@ -397,14 +412,9 @@ __kernel void m07800_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (left = final_len, off = 0; left >= 56; left -= 64, off += 16)
|
||||
{
|
||||
swap_buffer (&final[off]);
|
||||
|
||||
sha1_transform (&final[off + 0], &final[off + 4], &final[off + 8], &final[off + 12], digest);
|
||||
}
|
||||
|
||||
swap_buffer (&final[off]);
|
||||
|
||||
final[off + 14] = 0;
|
||||
final[off + 15] = final_len * 8;
|
||||
|
||||
sha1_transform (&final[off + 0], &final[off + 4], &final[off + 8], &final[off + 12], digest);
|
||||
@ -545,6 +555,22 @@ __kernel void m07800_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
final[13] = swap32 (w3[1] | s3[1]);
|
||||
final[14] = 0;
|
||||
final[15] = pw_salt_len * 8;
|
||||
final[16] = 0;
|
||||
final[17] = 0;
|
||||
final[18] = 0;
|
||||
final[19] = 0;
|
||||
final[20] = 0;
|
||||
final[21] = 0;
|
||||
final[22] = 0;
|
||||
final[23] = 0;
|
||||
final[24] = 0;
|
||||
final[25] = 0;
|
||||
final[26] = 0;
|
||||
final[27] = 0;
|
||||
final[28] = 0;
|
||||
final[29] = 0;
|
||||
final[30] = 0;
|
||||
final[31] = 0;
|
||||
|
||||
u32 digest[5];
|
||||
|
||||
@ -590,21 +616,24 @@ __kernel void m07800_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
digest[3] = SHA1M_D;
|
||||
digest[4] = SHA1M_E;
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (int i = 0; i < 32; i++) final[i] = 0;
|
||||
final[ 0] = swap32_S (w0[0]);
|
||||
final[ 1] = swap32_S (w0[1]);
|
||||
final[ 2] = swap32_S (w0[2]);
|
||||
final[ 3] = swap32_S (w0[3]);
|
||||
final[ 4] = swap32_S (w1[0]);
|
||||
final[ 5] = swap32_S (w1[1]);
|
||||
final[ 6] = swap32_S (w1[2]);
|
||||
final[ 7] = swap32_S (w1[3]);
|
||||
final[ 8] = 0;
|
||||
final[ 9] = 0;
|
||||
final[10] = 0;
|
||||
final[11] = 0;
|
||||
final[12] = 0;
|
||||
final[13] = 0;
|
||||
final[14] = 0;
|
||||
final[15] = 0;
|
||||
|
||||
final[0] = w0[0];
|
||||
final[1] = w0[1];
|
||||
final[2] = w0[2];
|
||||
final[3] = w0[3];
|
||||
final[4] = w1[0];
|
||||
final[5] = w1[1];
|
||||
final[6] = w1[2];
|
||||
final[7] = w1[3];
|
||||
|
||||
u32 final_len = out_len;
|
||||
u32 final_len = pw_len;
|
||||
|
||||
int i;
|
||||
|
||||
@ -612,14 +641,14 @@ __kernel void m07800_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (i = 0; i < lengthMagicArray - 4; i += 4)
|
||||
{
|
||||
const u32 tmp = GETSHIFTEDINT (theMagicArray, offsetMagicArray + i);
|
||||
const u32 tmp = GETSHIFTEDINT_CONST (theMagicArray, offsetMagicArray + i);
|
||||
|
||||
SETSHIFTEDINT (final, final_len + i, tmp);
|
||||
}
|
||||
|
||||
const u32 mask = 0xffffffff >> (((i - lengthMagicArray) & 3) * 8);
|
||||
const u32 mask = 0xffffffff << (((4 - (lengthMagicArray - i)) & 3) * 8);
|
||||
|
||||
const u32 tmp = GETSHIFTEDINT (theMagicArray, offsetMagicArray + i) & mask;
|
||||
const u32 tmp = GETSHIFTEDINT_CONST (theMagicArray, offsetMagicArray + i) & mask;
|
||||
|
||||
SETSHIFTEDINT (final, final_len + i, tmp);
|
||||
|
||||
@ -629,7 +658,7 @@ __kernel void m07800_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (i = 0; i < salt_len + 1; i += 4) // +1 for the 0x80
|
||||
{
|
||||
const u32 tmp = salt_buf[i / 4]; // attention, int[] not char[]
|
||||
const u32 tmp = swap32_S (salt_buf[i / 4]); // attention, int[] not char[]
|
||||
|
||||
SETSHIFTEDINT (final, final_len + i, tmp);
|
||||
}
|
||||
@ -643,14 +672,9 @@ __kernel void m07800_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (left = final_len, off = 0; left >= 56; left -= 64, off += 16)
|
||||
{
|
||||
swap_buffer (&final[off]);
|
||||
|
||||
sha1_transform (&final[off + 0], &final[off + 4], &final[off + 8], &final[off + 12], digest);
|
||||
}
|
||||
|
||||
swap_buffer (&final[off]);
|
||||
|
||||
final[off + 14] = 0;
|
||||
final[off + 15] = final_len * 8;
|
||||
|
||||
sha1_transform (&final[off + 0], &final[off + 4], &final[off + 8], &final[off + 12], digest);
|
||||
|
||||
@ -15,46 +15,42 @@
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
|
||||
#define GETSHIFTEDINT(a,n) amd_bytealign ((a)[((n)/4)+1], (a)[((n)/4)+0], (n))
|
||||
|
||||
#define SETSHIFTEDINT(a,n,v) \
|
||||
{ \
|
||||
const u32 s = ((n) & 3) * 8; \
|
||||
const u64 x = (u64) (v) << s; \
|
||||
(a)[((n)/4)+0] |= x; \
|
||||
(a)[((n)/4)+1] = x >> 32; \
|
||||
}
|
||||
|
||||
__constant u32 theMagicArray[64] =
|
||||
{
|
||||
0x1451ac91,0x4354679f,0xe03be724,0xc27b7428,0xeb133386,0x5ccb4f5a,0x37730a08,0x2f1c5d0e,
|
||||
0xe5e68f33,0xddae9bf8,0x8d4bf216,0xdcd4e12c,0x9ddfcbb0,0x176d70d4,0x3f424df9,0x94111b9b,
|
||||
0x9bc15b9f,0x039d0506,0x8a135e9d,0xe86a9a1e,0x17147cd9,0xf62ac758,0x0a6399a1,0xc370fdd7,
|
||||
0x13745ef6,0x040bc903,0x26f79826,0x2593928a,0x230da2b0,0x6d7963ed,0x3cfa3213,0xa39a0235,
|
||||
0x0a8eddb3,0xc351bf24,0x9f55cd7c,0x4c94af37,0x82520829,0x374e3bb2,0x9107179f,0xcdfd3b11,
|
||||
0, 0, 0, 0, 0, 0, 0, 0,
|
||||
0, 0, 0, 0, 0, 0, 0, 0,
|
||||
0, 0, 0, 0, 0, 0, 0, 0
|
||||
0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f,
|
||||
0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194,
|
||||
0x9f5bc19b, 0x06059d03, 0x9d5e138a, 0x1e9a6ae8, 0xd97c1417, 0x58c72af6, 0xa199630a, 0xd7fd70c3,
|
||||
0xf65e7413, 0x03c90b04, 0x2698f726, 0x8a929325, 0xb0a20d23, 0xed63796d, 0x1332fa3c, 0x35029aa3,
|
||||
0xb3dd8e0a, 0x24bf51c3, 0x7ccd559f, 0x37af944c, 0x29085282, 0xb23b4e37, 0x9f170791, 0x113bfdcd,
|
||||
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
|
||||
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
|
||||
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
|
||||
};
|
||||
|
||||
void swap_buffer (u32 final[16])
|
||||
static u32 GETSHIFTEDINT_CONST (__constant u32 *a, const int n)
|
||||
{
|
||||
final[ 0] = swap32 (final[ 0]);
|
||||
final[ 1] = swap32 (final[ 1]);
|
||||
final[ 2] = swap32 (final[ 2]);
|
||||
final[ 3] = swap32 (final[ 3]);
|
||||
final[ 4] = swap32 (final[ 4]);
|
||||
final[ 5] = swap32 (final[ 5]);
|
||||
final[ 6] = swap32 (final[ 6]);
|
||||
final[ 7] = swap32 (final[ 7]);
|
||||
final[ 8] = swap32 (final[ 8]);
|
||||
final[ 9] = swap32 (final[ 9]);
|
||||
final[10] = swap32 (final[10]);
|
||||
final[11] = swap32 (final[11]);
|
||||
final[12] = swap32 (final[12]);
|
||||
final[13] = swap32 (final[13]);
|
||||
final[14] = swap32 (final[14]);
|
||||
final[15] = swap32 (final[15]);
|
||||
const int d = n / 4;
|
||||
const int m = n & 3;
|
||||
|
||||
u64 tmp = hl32_to_64_S (a[d + 0], a[d + 1]);
|
||||
|
||||
tmp <<= m * 8;
|
||||
|
||||
return h32_from_64_S (tmp);
|
||||
}
|
||||
|
||||
|
||||
static void SETSHIFTEDINT (u32 *a, const int n, const u32 v)
|
||||
{
|
||||
const int d = n / 4;
|
||||
const int m = n & 3;
|
||||
|
||||
u64 tmp = hl32_to_64_S (v, 0);
|
||||
|
||||
tmp >>= m * 8;
|
||||
|
||||
a[d + 0] |= h32_from_64_S (tmp);
|
||||
a[d + 1] = l32_from_64_S (tmp);
|
||||
}
|
||||
|
||||
void sha1_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[5])
|
||||
@ -357,6 +353,22 @@ __kernel void m07800_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
final[13] = swap32 (w3[1] | s3[1]);
|
||||
final[14] = 0;
|
||||
final[15] = pw_salt_len * 8;
|
||||
final[16] = 0;
|
||||
final[17] = 0;
|
||||
final[18] = 0;
|
||||
final[19] = 0;
|
||||
final[20] = 0;
|
||||
final[21] = 0;
|
||||
final[22] = 0;
|
||||
final[23] = 0;
|
||||
final[24] = 0;
|
||||
final[25] = 0;
|
||||
final[26] = 0;
|
||||
final[27] = 0;
|
||||
final[28] = 0;
|
||||
final[29] = 0;
|
||||
final[30] = 0;
|
||||
final[31] = 0;
|
||||
|
||||
u32 digest[5];
|
||||
|
||||
@ -402,19 +414,22 @@ __kernel void m07800_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
digest[3] = SHA1M_D;
|
||||
digest[4] = SHA1M_E;
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (int i = 0; i < 32; i++) final[i] = 0;
|
||||
|
||||
final[0] = w0[0];
|
||||
final[1] = w0[1];
|
||||
final[2] = w0[2];
|
||||
final[3] = w0[3];
|
||||
final[4] = w1[0];
|
||||
final[5] = w1[1];
|
||||
final[6] = w1[2];
|
||||
final[7] = w1[3];
|
||||
final[ 0] = swap32_S (w0[0]);
|
||||
final[ 1] = swap32_S (w0[1]);
|
||||
final[ 2] = swap32_S (w0[2]);
|
||||
final[ 3] = swap32_S (w0[3]);
|
||||
final[ 4] = swap32_S (w1[0]);
|
||||
final[ 5] = swap32_S (w1[1]);
|
||||
final[ 6] = swap32_S (w1[2]);
|
||||
final[ 7] = swap32_S (w1[3]);
|
||||
final[ 8] = swap32_S (w2[0]);
|
||||
final[ 9] = swap32_S (w2[1]);
|
||||
final[10] = swap32_S (w2[2]);
|
||||
final[11] = swap32_S (w2[3]);
|
||||
final[12] = swap32_S (w3[0]);
|
||||
final[13] = swap32_S (w3[1]);
|
||||
final[14] = 0;
|
||||
final[15] = 0;
|
||||
|
||||
u32 final_len = pw_len;
|
||||
|
||||
@ -424,14 +439,14 @@ __kernel void m07800_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (i = 0; i < lengthMagicArray - 4; i += 4)
|
||||
{
|
||||
const u32 tmp = GETSHIFTEDINT (theMagicArray, offsetMagicArray + i);
|
||||
const u32 tmp = GETSHIFTEDINT_CONST (theMagicArray, offsetMagicArray + i);
|
||||
|
||||
SETSHIFTEDINT (final, final_len + i, tmp);
|
||||
}
|
||||
|
||||
const u32 mask = 0xffffffff >> (((i - lengthMagicArray) & 3) * 8);
|
||||
const u32 mask = 0xffffffff << (((4 - (lengthMagicArray - i)) & 3) * 8);
|
||||
|
||||
const u32 tmp = GETSHIFTEDINT (theMagicArray, offsetMagicArray + i) & mask;
|
||||
const u32 tmp = GETSHIFTEDINT_CONST (theMagicArray, offsetMagicArray + i) & mask;
|
||||
|
||||
SETSHIFTEDINT (final, final_len + i, tmp);
|
||||
|
||||
@ -441,7 +456,7 @@ __kernel void m07800_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (i = 0; i < salt_len + 1; i += 4) // +1 for the 0x80
|
||||
{
|
||||
const u32 tmp = salt_buf[i / 4]; // attention, int[] not char[]
|
||||
const u32 tmp = swap32_S (salt_buf[i / 4]); // attention, int[] not char[]
|
||||
|
||||
SETSHIFTEDINT (final, final_len + i, tmp);
|
||||
}
|
||||
@ -455,14 +470,9 @@ __kernel void m07800_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (left = final_len, off = 0; left >= 56; left -= 64, off += 16)
|
||||
{
|
||||
swap_buffer (&final[off]);
|
||||
|
||||
sha1_transform (&final[off + 0], &final[off + 4], &final[off + 8], &final[off + 12], digest);
|
||||
}
|
||||
|
||||
swap_buffer (&final[off]);
|
||||
|
||||
final[off + 14] = 0;
|
||||
final[off + 15] = final_len * 8;
|
||||
|
||||
sha1_transform (&final[off + 0], &final[off + 4], &final[off + 8], &final[off + 12], digest);
|
||||
@ -663,6 +673,22 @@ __kernel void m07800_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
final[13] = swap32 (w3[1] | s3[1]);
|
||||
final[14] = 0;
|
||||
final[15] = pw_salt_len * 8;
|
||||
final[16] = 0;
|
||||
final[17] = 0;
|
||||
final[18] = 0;
|
||||
final[19] = 0;
|
||||
final[20] = 0;
|
||||
final[21] = 0;
|
||||
final[22] = 0;
|
||||
final[23] = 0;
|
||||
final[24] = 0;
|
||||
final[25] = 0;
|
||||
final[26] = 0;
|
||||
final[27] = 0;
|
||||
final[28] = 0;
|
||||
final[29] = 0;
|
||||
final[30] = 0;
|
||||
final[31] = 0;
|
||||
|
||||
u32 digest[5];
|
||||
|
||||
@ -708,19 +734,22 @@ __kernel void m07800_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
digest[3] = SHA1M_D;
|
||||
digest[4] = SHA1M_E;
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (int i = 0; i < 32; i++) final[i] = 0;
|
||||
|
||||
final[0] = w0[0];
|
||||
final[1] = w0[1];
|
||||
final[2] = w0[2];
|
||||
final[3] = w0[3];
|
||||
final[4] = w1[0];
|
||||
final[5] = w1[1];
|
||||
final[6] = w1[2];
|
||||
final[7] = w1[3];
|
||||
final[ 0] = swap32_S (w0[0]);
|
||||
final[ 1] = swap32_S (w0[1]);
|
||||
final[ 2] = swap32_S (w0[2]);
|
||||
final[ 3] = swap32_S (w0[3]);
|
||||
final[ 4] = swap32_S (w1[0]);
|
||||
final[ 5] = swap32_S (w1[1]);
|
||||
final[ 6] = swap32_S (w1[2]);
|
||||
final[ 7] = swap32_S (w1[3]);
|
||||
final[ 8] = swap32_S (w2[0]);
|
||||
final[ 9] = swap32_S (w2[1]);
|
||||
final[10] = swap32_S (w2[2]);
|
||||
final[11] = swap32_S (w2[3]);
|
||||
final[12] = swap32_S (w3[0]);
|
||||
final[13] = swap32_S (w3[1]);
|
||||
final[14] = 0;
|
||||
final[15] = 0;
|
||||
|
||||
u32 final_len = pw_len;
|
||||
|
||||
@ -730,14 +759,14 @@ __kernel void m07800_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (i = 0; i < lengthMagicArray - 4; i += 4)
|
||||
{
|
||||
const u32 tmp = GETSHIFTEDINT (theMagicArray, offsetMagicArray + i);
|
||||
const u32 tmp = GETSHIFTEDINT_CONST (theMagicArray, offsetMagicArray + i);
|
||||
|
||||
SETSHIFTEDINT (final, final_len + i, tmp);
|
||||
}
|
||||
|
||||
const u32 mask = 0xffffffff >> (((i - lengthMagicArray) & 3) * 8);
|
||||
const u32 mask = 0xffffffff << (((4 - (lengthMagicArray - i)) & 3) * 8);
|
||||
|
||||
const u32 tmp = GETSHIFTEDINT (theMagicArray, offsetMagicArray + i) & mask;
|
||||
const u32 tmp = GETSHIFTEDINT_CONST (theMagicArray, offsetMagicArray + i) & mask;
|
||||
|
||||
SETSHIFTEDINT (final, final_len + i, tmp);
|
||||
|
||||
@ -747,7 +776,7 @@ __kernel void m07800_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (i = 0; i < salt_len + 1; i += 4) // +1 for the 0x80
|
||||
{
|
||||
const u32 tmp = salt_buf[i / 4]; // attention, int[] not char[]
|
||||
const u32 tmp = swap32_S (salt_buf[i / 4]); // attention, int[] not char[]
|
||||
|
||||
SETSHIFTEDINT (final, final_len + i, tmp);
|
||||
}
|
||||
@ -761,14 +790,9 @@ __kernel void m07800_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (left = final_len, off = 0; left >= 56; left -= 64, off += 16)
|
||||
{
|
||||
swap_buffer (&final[off]);
|
||||
|
||||
sha1_transform (&final[off + 0], &final[off + 4], &final[off + 8], &final[off + 12], digest);
|
||||
}
|
||||
|
||||
swap_buffer (&final[off]);
|
||||
|
||||
final[off + 14] = 0;
|
||||
final[off + 15] = final_len * 8;
|
||||
|
||||
sha1_transform (&final[off + 0], &final[off + 4], &final[off + 8], &final[off + 12], digest);
|
||||
|
||||
@ -15,46 +15,42 @@
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
|
||||
#define GETSHIFTEDINT(a,n) amd_bytealign ((a)[((n)/4)+1], (a)[((n)/4)+0], (n))
|
||||
|
||||
#define SETSHIFTEDINT(a,n,v) \
|
||||
{ \
|
||||
const u32 s = ((n) & 3) * 8; \
|
||||
const u64 x = (u64) (v) << s; \
|
||||
(a)[((n)/4)+0] |= x; \
|
||||
(a)[((n)/4)+1] = x >> 32; \
|
||||
}
|
||||
|
||||
__constant u32 theMagicArray[64] =
|
||||
{
|
||||
0x1451ac91,0x4354679f,0xe03be724,0xc27b7428,0xeb133386,0x5ccb4f5a,0x37730a08,0x2f1c5d0e,
|
||||
0xe5e68f33,0xddae9bf8,0x8d4bf216,0xdcd4e12c,0x9ddfcbb0,0x176d70d4,0x3f424df9,0x94111b9b,
|
||||
0x9bc15b9f,0x039d0506,0x8a135e9d,0xe86a9a1e,0x17147cd9,0xf62ac758,0x0a6399a1,0xc370fdd7,
|
||||
0x13745ef6,0x040bc903,0x26f79826,0x2593928a,0x230da2b0,0x6d7963ed,0x3cfa3213,0xa39a0235,
|
||||
0x0a8eddb3,0xc351bf24,0x9f55cd7c,0x4c94af37,0x82520829,0x374e3bb2,0x9107179f,0xcdfd3b11,
|
||||
0, 0, 0, 0, 0, 0, 0, 0,
|
||||
0, 0, 0, 0, 0, 0, 0, 0,
|
||||
0, 0, 0, 0, 0, 0, 0, 0
|
||||
0x91ac5114, 0x9f675443, 0x24e73be0, 0x28747bc2, 0x863313eb, 0x5a4fcb5c, 0x080a7337, 0x0e5d1c2f,
|
||||
0x338fe6e5, 0xf89baedd, 0x16f24b8d, 0x2ce1d4dc, 0xb0cbdf9d, 0xd4706d17, 0xf94d423f, 0x9b1b1194,
|
||||
0x9f5bc19b, 0x06059d03, 0x9d5e138a, 0x1e9a6ae8, 0xd97c1417, 0x58c72af6, 0xa199630a, 0xd7fd70c3,
|
||||
0xf65e7413, 0x03c90b04, 0x2698f726, 0x8a929325, 0xb0a20d23, 0xed63796d, 0x1332fa3c, 0x35029aa3,
|
||||
0xb3dd8e0a, 0x24bf51c3, 0x7ccd559f, 0x37af944c, 0x29085282, 0xb23b4e37, 0x9f170791, 0x113bfdcd,
|
||||
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
|
||||
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
|
||||
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
|
||||
};
|
||||
|
||||
void swap_buffer (u32 final[16])
|
||||
static u32 GETSHIFTEDINT_CONST (__constant u32 *a, const int n)
|
||||
{
|
||||
final[ 0] = swap32 (final[ 0]);
|
||||
final[ 1] = swap32 (final[ 1]);
|
||||
final[ 2] = swap32 (final[ 2]);
|
||||
final[ 3] = swap32 (final[ 3]);
|
||||
final[ 4] = swap32 (final[ 4]);
|
||||
final[ 5] = swap32 (final[ 5]);
|
||||
final[ 6] = swap32 (final[ 6]);
|
||||
final[ 7] = swap32 (final[ 7]);
|
||||
final[ 8] = swap32 (final[ 8]);
|
||||
final[ 9] = swap32 (final[ 9]);
|
||||
final[10] = swap32 (final[10]);
|
||||
final[11] = swap32 (final[11]);
|
||||
final[12] = swap32 (final[12]);
|
||||
final[13] = swap32 (final[13]);
|
||||
final[14] = swap32 (final[14]);
|
||||
final[15] = swap32 (final[15]);
|
||||
const int d = n / 4;
|
||||
const int m = n & 3;
|
||||
|
||||
u64 tmp = hl32_to_64_S (a[d + 0], a[d + 1]);
|
||||
|
||||
tmp <<= m * 8;
|
||||
|
||||
return h32_from_64_S (tmp);
|
||||
}
|
||||
|
||||
|
||||
static void SETSHIFTEDINT (u32 *a, const int n, const u32 v)
|
||||
{
|
||||
const int d = n / 4;
|
||||
const int m = n & 3;
|
||||
|
||||
u64 tmp = hl32_to_64_S (v, 0);
|
||||
|
||||
tmp >>= m * 8;
|
||||
|
||||
a[d + 0] |= h32_from_64_S (tmp);
|
||||
a[d + 1] = l32_from_64_S (tmp);
|
||||
}
|
||||
|
||||
void sha1_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[5])
|
||||
@ -200,14 +196,14 @@ void m07800m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
|
||||
|
||||
u32 salt_buf[8];
|
||||
|
||||
salt_buf[0] = salt_bufs[salt_pos].salt_buf[0];
|
||||
salt_buf[1] = salt_bufs[salt_pos].salt_buf[1];
|
||||
salt_buf[2] = salt_bufs[salt_pos].salt_buf[2];
|
||||
salt_buf[3] = salt_bufs[salt_pos].salt_buf[3];
|
||||
salt_buf[4] = salt_bufs[salt_pos].salt_buf[4];
|
||||
salt_buf[5] = salt_bufs[salt_pos].salt_buf[5];
|
||||
salt_buf[6] = salt_bufs[salt_pos].salt_buf[6];
|
||||
salt_buf[7] = salt_bufs[salt_pos].salt_buf[7];
|
||||
salt_buf[0] = swap32_S (salt_bufs[salt_pos].salt_buf[0]);
|
||||
salt_buf[1] = swap32_S (salt_bufs[salt_pos].salt_buf[1]);
|
||||
salt_buf[2] = swap32_S (salt_bufs[salt_pos].salt_buf[2]);
|
||||
salt_buf[3] = swap32_S (salt_bufs[salt_pos].salt_buf[3]);
|
||||
salt_buf[4] = swap32_S (salt_bufs[salt_pos].salt_buf[4]);
|
||||
salt_buf[5] = swap32_S (salt_bufs[salt_pos].salt_buf[5]);
|
||||
salt_buf[6] = swap32_S (salt_bufs[salt_pos].salt_buf[6]);
|
||||
salt_buf[7] = swap32_S (salt_bufs[salt_pos].salt_buf[7]);
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -216,14 +212,14 @@ void m07800m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
|
||||
u32 s2[4];
|
||||
u32 s3[4];
|
||||
|
||||
s0[0] = swap32 (salt_buf[0]);
|
||||
s0[1] = swap32 (salt_buf[1]);
|
||||
s0[2] = swap32 (salt_buf[2]);
|
||||
s0[3] = swap32 (salt_buf[3]);
|
||||
s1[0] = swap32 (salt_buf[4]);
|
||||
s1[1] = swap32 (salt_buf[5]);
|
||||
s1[2] = swap32 (salt_buf[6]);
|
||||
s1[3] = swap32 (salt_buf[7]);
|
||||
s0[0] = salt_buf[0];
|
||||
s0[1] = salt_buf[1];
|
||||
s0[2] = salt_buf[2];
|
||||
s0[3] = salt_buf[3];
|
||||
s1[0] = salt_buf[4];
|
||||
s1[1] = salt_buf[5];
|
||||
s1[2] = salt_buf[6];
|
||||
s1[3] = salt_buf[7];
|
||||
s2[0] = 0;
|
||||
s2[1] = 0;
|
||||
s2[2] = 0;
|
||||
@ -273,6 +269,22 @@ void m07800m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
|
||||
final[13] = w3[1] | s3[1];
|
||||
final[14] = 0;
|
||||
final[15] = pw_salt_len * 8;
|
||||
final[16] = 0;
|
||||
final[17] = 0;
|
||||
final[18] = 0;
|
||||
final[19] = 0;
|
||||
final[20] = 0;
|
||||
final[21] = 0;
|
||||
final[22] = 0;
|
||||
final[23] = 0;
|
||||
final[24] = 0;
|
||||
final[25] = 0;
|
||||
final[26] = 0;
|
||||
final[27] = 0;
|
||||
final[28] = 0;
|
||||
final[29] = 0;
|
||||
final[30] = 0;
|
||||
final[31] = 0;
|
||||
|
||||
u32 digest[5];
|
||||
|
||||
@ -318,19 +330,22 @@ void m07800m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
|
||||
digest[3] = SHA1M_D;
|
||||
digest[4] = SHA1M_E;
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (int i = 0; i < 32; i++) final[i] = 0;
|
||||
|
||||
final[0] = swap32 (w0[0]);
|
||||
final[1] = swap32 (w0[1]);
|
||||
final[2] = swap32 (w0[2]);
|
||||
final[3] = swap32 (w0[3]);
|
||||
final[4] = swap32 (w1[0]);
|
||||
final[5] = swap32 (w1[1]);
|
||||
final[6] = swap32 (w1[2]);
|
||||
final[7] = swap32 (w1[3]);
|
||||
final[ 0] = w0[0];
|
||||
final[ 1] = w0[1];
|
||||
final[ 2] = w0[2];
|
||||
final[ 3] = w0[3];
|
||||
final[ 4] = w1[0];
|
||||
final[ 5] = w1[1];
|
||||
final[ 6] = w1[2];
|
||||
final[ 7] = w1[3];
|
||||
final[ 8] = 0;
|
||||
final[ 9] = 0;
|
||||
final[10] = 0;
|
||||
final[11] = 0;
|
||||
final[12] = 0;
|
||||
final[13] = 0;
|
||||
final[14] = 0;
|
||||
final[15] = 0;
|
||||
|
||||
u32 final_len = pw_len;
|
||||
|
||||
@ -340,14 +355,14 @@ void m07800m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
|
||||
|
||||
for (i = 0; i < lengthMagicArray - 4; i += 4)
|
||||
{
|
||||
const u32 tmp = GETSHIFTEDINT (theMagicArray, offsetMagicArray + i);
|
||||
const u32 tmp = GETSHIFTEDINT_CONST (theMagicArray, offsetMagicArray + i);
|
||||
|
||||
SETSHIFTEDINT (final, final_len + i, tmp);
|
||||
}
|
||||
|
||||
const u32 mask = 0xffffffff >> (((i - lengthMagicArray) & 3) * 8);
|
||||
const u32 mask = 0xffffffff << (((4 - (lengthMagicArray - i)) & 3) * 8);
|
||||
|
||||
const u32 tmp = GETSHIFTEDINT (theMagicArray, offsetMagicArray + i) & mask;
|
||||
const u32 tmp = GETSHIFTEDINT_CONST (theMagicArray, offsetMagicArray + i) & mask;
|
||||
|
||||
SETSHIFTEDINT (final, final_len + i, tmp);
|
||||
|
||||
@ -371,14 +386,9 @@ void m07800m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
|
||||
|
||||
for (left = final_len, off = 0; left >= 56; left -= 64, off += 16)
|
||||
{
|
||||
swap_buffer (&final[off]);
|
||||
|
||||
sha1_transform (&final[off + 0], &final[off + 4], &final[off + 8], &final[off + 12], digest);
|
||||
}
|
||||
|
||||
swap_buffer (&final[off]);
|
||||
|
||||
final[off + 14] = 0;
|
||||
final[off + 15] = final_len * 8;
|
||||
|
||||
sha1_transform (&final[off + 0], &final[off + 4], &final[off + 8], &final[off + 12], digest);
|
||||
@ -402,14 +412,14 @@ void m07800s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
|
||||
|
||||
u32 salt_buf[8];
|
||||
|
||||
salt_buf[0] = salt_bufs[salt_pos].salt_buf[0];
|
||||
salt_buf[1] = salt_bufs[salt_pos].salt_buf[1];
|
||||
salt_buf[2] = salt_bufs[salt_pos].salt_buf[2];
|
||||
salt_buf[3] = salt_bufs[salt_pos].salt_buf[3];
|
||||
salt_buf[4] = salt_bufs[salt_pos].salt_buf[4];
|
||||
salt_buf[5] = salt_bufs[salt_pos].salt_buf[5];
|
||||
salt_buf[6] = salt_bufs[salt_pos].salt_buf[6];
|
||||
salt_buf[7] = salt_bufs[salt_pos].salt_buf[7];
|
||||
salt_buf[0] = swap32_S (salt_bufs[salt_pos].salt_buf[0]);
|
||||
salt_buf[1] = swap32_S (salt_bufs[salt_pos].salt_buf[1]);
|
||||
salt_buf[2] = swap32_S (salt_bufs[salt_pos].salt_buf[2]);
|
||||
salt_buf[3] = swap32_S (salt_bufs[salt_pos].salt_buf[3]);
|
||||
salt_buf[4] = swap32_S (salt_bufs[salt_pos].salt_buf[4]);
|
||||
salt_buf[5] = swap32_S (salt_bufs[salt_pos].salt_buf[5]);
|
||||
salt_buf[6] = swap32_S (salt_bufs[salt_pos].salt_buf[6]);
|
||||
salt_buf[7] = swap32_S (salt_bufs[salt_pos].salt_buf[7]);
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -418,14 +428,14 @@ void m07800s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
|
||||
u32 s2[4];
|
||||
u32 s3[4];
|
||||
|
||||
s0[0] = swap32 (salt_buf[0]);
|
||||
s0[1] = swap32 (salt_buf[1]);
|
||||
s0[2] = swap32 (salt_buf[2]);
|
||||
s0[3] = swap32 (salt_buf[3]);
|
||||
s1[0] = swap32 (salt_buf[4]);
|
||||
s1[1] = swap32 (salt_buf[5]);
|
||||
s1[2] = swap32 (salt_buf[6]);
|
||||
s1[3] = swap32 (salt_buf[7]);
|
||||
s0[0] = salt_buf[0];
|
||||
s0[1] = salt_buf[1];
|
||||
s0[2] = salt_buf[2];
|
||||
s0[3] = salt_buf[3];
|
||||
s1[0] = salt_buf[4];
|
||||
s1[1] = salt_buf[5];
|
||||
s1[2] = salt_buf[6];
|
||||
s1[3] = salt_buf[7];
|
||||
s2[0] = 0;
|
||||
s2[1] = 0;
|
||||
s2[2] = 0;
|
||||
@ -487,6 +497,22 @@ void m07800s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
|
||||
final[13] = w3[1] | s3[1];
|
||||
final[14] = 0;
|
||||
final[15] = pw_salt_len * 8;
|
||||
final[16] = 0;
|
||||
final[17] = 0;
|
||||
final[18] = 0;
|
||||
final[19] = 0;
|
||||
final[20] = 0;
|
||||
final[21] = 0;
|
||||
final[22] = 0;
|
||||
final[23] = 0;
|
||||
final[24] = 0;
|
||||
final[25] = 0;
|
||||
final[26] = 0;
|
||||
final[27] = 0;
|
||||
final[28] = 0;
|
||||
final[29] = 0;
|
||||
final[30] = 0;
|
||||
final[31] = 0;
|
||||
|
||||
u32 digest[5];
|
||||
|
||||
@ -532,19 +558,22 @@ void m07800s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
|
||||
digest[3] = SHA1M_D;
|
||||
digest[4] = SHA1M_E;
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (int i = 0; i < 32; i++) final[i] = 0;
|
||||
|
||||
final[0] = swap32 (w0[0]);
|
||||
final[1] = swap32 (w0[1]);
|
||||
final[2] = swap32 (w0[2]);
|
||||
final[3] = swap32 (w0[3]);
|
||||
final[4] = swap32 (w1[0]);
|
||||
final[5] = swap32 (w1[1]);
|
||||
final[6] = swap32 (w1[2]);
|
||||
final[7] = swap32 (w1[3]);
|
||||
final[ 0] = w0[0];
|
||||
final[ 1] = w0[1];
|
||||
final[ 2] = w0[2];
|
||||
final[ 3] = w0[3];
|
||||
final[ 4] = w1[0];
|
||||
final[ 5] = w1[1];
|
||||
final[ 6] = w1[2];
|
||||
final[ 7] = w1[3];
|
||||
final[ 8] = 0;
|
||||
final[ 9] = 0;
|
||||
final[10] = 0;
|
||||
final[11] = 0;
|
||||
final[12] = 0;
|
||||
final[13] = 0;
|
||||
final[14] = 0;
|
||||
final[15] = 0;
|
||||
|
||||
u32 final_len = pw_len;
|
||||
|
||||
@ -554,14 +583,14 @@ void m07800s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
|
||||
|
||||
for (i = 0; i < lengthMagicArray - 4; i += 4)
|
||||
{
|
||||
const u32 tmp = GETSHIFTEDINT (theMagicArray, offsetMagicArray + i);
|
||||
const u32 tmp = GETSHIFTEDINT_CONST (theMagicArray, offsetMagicArray + i);
|
||||
|
||||
SETSHIFTEDINT (final, final_len + i, tmp);
|
||||
}
|
||||
|
||||
const u32 mask = 0xffffffff >> (((i - lengthMagicArray) & 3) * 8);
|
||||
const u32 mask = 0xffffffff << (((4 - (lengthMagicArray - i)) & 3) * 8);
|
||||
|
||||
const u32 tmp = GETSHIFTEDINT (theMagicArray, offsetMagicArray + i) & mask;
|
||||
const u32 tmp = GETSHIFTEDINT_CONST (theMagicArray, offsetMagicArray + i) & mask;
|
||||
|
||||
SETSHIFTEDINT (final, final_len + i, tmp);
|
||||
|
||||
@ -585,14 +614,9 @@ void m07800s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
|
||||
|
||||
for (left = final_len, off = 0; left >= 56; left -= 64, off += 16)
|
||||
{
|
||||
swap_buffer (&final[off]);
|
||||
|
||||
sha1_transform (&final[off + 0], &final[off + 4], &final[off + 8], &final[off + 12], digest);
|
||||
}
|
||||
|
||||
swap_buffer (&final[off]);
|
||||
|
||||
final[off + 14] = 0;
|
||||
final[off + 15] = final_len * 8;
|
||||
|
||||
sha1_transform (&final[off + 0], &final[off + 4], &final[off + 8], &final[off + 12], digest);
|
||||
|
||||
Loading…
Reference in New Issue
Block a user