arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-11-29 19:38:18 +00:00
Code Issues Releases Wiki Activity

Move 0x80 to hardcoded position for sha3-256 bit in order to allow ROCm compiler to use registers only

Browse Source
This commit is contained in:
jsteube 2017-08-18 16:22:25 +02:00
parent c8da13c3aa
commit d9c906e134
3 changed files with 18 additions and 90 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

36
OpenCL/m05000_a0-optimized.cl
View File

@ -98,28 +98,18 @@ __kernel void m05000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
* constants * constants
*/ */
const u8 keccakf_rotc[24] = const int keccakf_rotc[24] =
{ {
1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14, 1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44 27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
}; };
const u8 keccakf_piln[24] = const int keccakf_piln[24] =
{ {
10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4, 10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1 15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
}; };
/**
* 0x80 keccak, very special
*/
const u32 mdlen = salt_bufs[salt_pos].keccak_mdlen;
const u32 rsiz = 200 - (2 * mdlen);
const u32 add80w = (rsiz - 1) / 8;
/** /**
* loop * loop
*/ */
@ -157,7 +147,7 @@ __kernel void m05000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
st[13] = 0; st[13] = 0;
st[14] = 0; st[14] = 0;
st[15] = 0; st[15] = 0;
st[16] = 0; st[16] = 0x8000000000000000;
st[17] = 0; st[17] = 0;
st[18] = 0; st[18] = 0;
st[19] = 0; st[19] = 0;
@ -167,8 +157,6 @@ __kernel void m05000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
st[23] = 0; st[23] = 0;
st[24] = 0; st[24] = 0;
st[add80w] |= 0x8000000000000000;
int round; int round;
for (round = 0; round < KECCAK_ROUNDS; round++) for (round = 0; round < KECCAK_ROUNDS; round++)
@ -282,28 +270,18 @@ __kernel void m05000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
* constants * constants
*/ */
const u8 keccakf_rotc[24] = const int keccakf_rotc[24] =
{ {
1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14, 1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44 27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
}; };
const u8 keccakf_piln[24] = const int keccakf_piln[24] =
{ {
10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4, 10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1 15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
}; };
/**
* 0x80 keccak, very special
*/
const u32 mdlen = salt_bufs[salt_pos].keccak_mdlen;
const u32 rsiz = 200 - (2 * mdlen);
const u32 add80w = (rsiz - 1) / 8;
/** /**
* digest * digest
*/ */
@ -353,7 +331,7 @@ __kernel void m05000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
st[13] = 0; st[13] = 0;
st[14] = 0; st[14] = 0;
st[15] = 0; st[15] = 0;
st[16] = 0; st[16] = 0x8000000000000000;
st[17] = 0; st[17] = 0;
st[18] = 0; st[18] = 0;
st[19] = 0; st[19] = 0;
@ -363,8 +341,6 @@ __kernel void m05000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
st[23] = 0; st[23] = 0;
st[24] = 0; st[24] = 0;
st[add80w] |= 0x8000000000000000;
int round; int round;
for (round = 0; round < KECCAK_ROUNDS; round++) for (round = 0; round < KECCAK_ROUNDS; round++)

36
OpenCL/m05000_a1-optimized.cl
View File

@ -96,28 +96,18 @@ __kernel void m05000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
* constants * constants
*/ */
const u8 keccakf_rotc[24] = const int keccakf_rotc[24] =
{ {
1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14, 1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44 27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
}; };
const u8 keccakf_piln[24] = const int keccakf_piln[24] =
{ {
10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4, 10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1 15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
}; };
/**
* 0x80 keccak, very special
*/
const u32 mdlen = salt_bufs[salt_pos].keccak_mdlen;
const u32 rsiz = 200 - (2 * mdlen);
const u32 add80w = (rsiz - 1) / 8;
/** /**
* loop * loop
*/ */
@ -213,7 +203,7 @@ __kernel void m05000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
st[13] = 0; st[13] = 0;
st[14] = 0; st[14] = 0;
st[15] = 0; st[15] = 0;
st[16] = 0; st[16] = 0x8000000000000000;
st[17] = 0; st[17] = 0;
st[18] = 0; st[18] = 0;
st[19] = 0; st[19] = 0;
@ -223,8 +213,6 @@ __kernel void m05000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
st[23] = 0; st[23] = 0;
st[24] = 0; st[24] = 0;
st[add80w] |= 0x8000000000000000;
int round; int round;
for (round = 0; round < KECCAK_ROUNDS; round++) for (round = 0; round < KECCAK_ROUNDS; round++)
@ -338,28 +326,18 @@ __kernel void m05000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
* constants * constants
*/ */
const u8 keccakf_rotc[24] = const int keccakf_rotc[24] =
{ {
1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14, 1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44 27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
}; };
const u8 keccakf_piln[24] = const int keccakf_piln[24] =
{ {
10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4, 10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1 15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
}; };
/**
* 0x80 keccak, very special
*/
const u32 mdlen = salt_bufs[salt_pos].keccak_mdlen;
const u32 rsiz = 200 - (2 * mdlen);
const u32 add80w = (rsiz - 1) / 8;
/** /**
* digest * digest
*/ */
@ -467,7 +445,7 @@ __kernel void m05000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
st[13] = 0; st[13] = 0;
st[14] = 0; st[14] = 0;
st[15] = 0; st[15] = 0;
st[16] = 0; st[16] = 0x8000000000000000;
st[17] = 0; st[17] = 0;
st[18] = 0; st[18] = 0;
st[19] = 0; st[19] = 0;
@ -477,8 +455,6 @@ __kernel void m05000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
st[23] = 0; st[23] = 0;
st[24] = 0; st[24] = 0;
st[add80w] |= 0x8000000000000000;
int round; int round;
for (round = 0; round < KECCAK_ROUNDS; round++) for (round = 0; round < KECCAK_ROUNDS; round++)

36
OpenCL/m05000_a3-optimized.cl
View File

@ -75,28 +75,18 @@ void m05000m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
* constants * constants
*/ */
const u8 keccakf_rotc[24] = const int keccakf_rotc[24] =
{ {
1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14, 1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44 27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
}; };
const u8 keccakf_piln[24] = const int keccakf_piln[24] =
{ {
10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4, 10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1 15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
}; };
/**
* 0x80 keccak, very special
*/
const u32 mdlen = salt_bufs[salt_pos].keccak_mdlen;
const u32 rsiz = 200 - (2 * mdlen);
const u32 add80w = (rsiz - 1) / 8;
/** /**
* loop * loop
*/ */
@ -153,7 +143,7 @@ void m05000m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
st[13] = 0; st[13] = 0;
st[14] = 0; st[14] = 0;
st[15] = 0; st[15] = 0;
st[16] = 0; st[16] = 0x8000000000000000;
st[17] = 0; st[17] = 0;
st[18] = 0; st[18] = 0;
st[19] = 0; st[19] = 0;
@ -163,8 +153,6 @@ void m05000m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
st[23] = 0; st[23] = 0;
st[24] = 0; st[24] = 0;
st[add80w] |= 0x8000000000000000;
int round; int round;
for (round = 0; round < KECCAK_ROUNDS; round++) for (round = 0; round < KECCAK_ROUNDS; round++)
@ -238,28 +226,18 @@ void m05000s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
* constants * constants
*/ */
const u8 keccakf_rotc[24] = const int keccakf_rotc[24] =
{ {
1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14, 1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44 27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
}; };
const u8 keccakf_piln[24] = const int keccakf_piln[24] =
{ {
10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4, 10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1 15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
}; };
/**
* 0x80 keccak, very special
*/
const u32 mdlen = salt_bufs[salt_pos].keccak_mdlen;
const u32 rsiz = 200 - (2 * mdlen);
const u32 add80w = (rsiz - 1) / 8;
/** /**
* digest * digest
*/ */
@ -328,7 +306,7 @@ void m05000s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
st[13] = 0; st[13] = 0;
st[14] = 0; st[14] = 0;
st[15] = 0; st[15] = 0;
st[16] = 0; st[16] = 0x8000000000000000;
st[17] = 0; st[17] = 0;
st[18] = 0; st[18] = 0;
st[19] = 0; st[19] = 0;
@ -338,8 +316,6 @@ void m05000s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
st[23] = 0; st[23] = 0;
st[24] = 0; st[24] = 0;
st[add80w] |= 0x8000000000000000;
int round; int round;
for (round = 0; round < KECCAK_ROUNDS; round++) for (round = 0; round < KECCAK_ROUNDS; round++)