arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2025-01-18 11:40:58 +00:00
Code Issues Releases Wiki Activity

Fix -m 5000 performance on NV after change to new standard

Browse Source
This commit is contained in:
jsteube 2016-04-21 09:01:31 +02:00
parent 444a5f4d52
commit ad01afeeb8
3 changed files with 96 additions and 36 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

44
OpenCL/m05000_a0.cl
View File

@ -34,18 +34,6 @@ __constant u64 keccakf_rndc[24] =
0x8000000000008080, 0x0000000080000001, 0x8000000080008008 0x8000000000008080, 0x0000000080000001, 0x8000000080008008
}; };
__constant u8 keccakf_rotc[24] =
{
1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
};
__constant u8 keccakf_piln[24] =
{
10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
};
#ifndef KECCAK_ROUNDS #ifndef KECCAK_ROUNDS
#define KECCAK_ROUNDS 24 #define KECCAK_ROUNDS 24
#endif #endif
@ -114,6 +102,22 @@ __kernel void m05000_m04 (__global pw_t *pws, __global kernel_rule_t * rules_bu
const u32 pw_len = pws[gid].pw_len; const u32 pw_len = pws[gid].pw_len;
/**
* constants
*/
const u8 keccakf_rotc[24] =
{
1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
};
const u8 keccakf_piln[24] =
{
10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
};
/** /**
* 0x80 keccak, very special * 0x80 keccak, very special
*/ */
@ -282,6 +286,22 @@ __kernel void m05000_s04 (__global pw_t *pws, __global kernel_rule_t * rules_bu
const u32 pw_len = pws[gid].pw_len; const u32 pw_len = pws[gid].pw_len;
/**
* constants
*/
const u8 keccakf_rotc[24] =
{
1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
};
const u8 keccakf_piln[24] =
{
10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
};
/** /**
* 0x80 keccak, very special * 0x80 keccak, very special
*/ */

44
OpenCL/m05000_a1.cl
View File

@ -32,18 +32,6 @@ __constant u64 keccakf_rndc[24] =
0x8000000000008080, 0x0000000080000001, 0x8000000080008008 0x8000000000008080, 0x0000000080000001, 0x8000000080008008
}; };
__constant u8 keccakf_rotc[24] =
{
1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
};
__constant u8 keccakf_piln[24] =
{
10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
};
#ifndef KECCAK_ROUNDS #ifndef KECCAK_ROUNDS
#define KECCAK_ROUNDS 24 #define KECCAK_ROUNDS 24
#endif #endif
@ -112,6 +100,22 @@ __kernel void m05000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
const u32 pw_l_len = pws[gid].pw_len; const u32 pw_l_len = pws[gid].pw_len;
/**
* constants
*/
const u8 keccakf_rotc[24] =
{
1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
};
const u8 keccakf_piln[24] =
{
10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
};
/** /**
* 0x80 keccak, very special * 0x80 keccak, very special
*/ */
@ -338,6 +342,22 @@ __kernel void m05000_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
const u32 pw_l_len = pws[gid].pw_len; const u32 pw_l_len = pws[gid].pw_len;
/**
* constants
*/
const u8 keccakf_rotc[24] =
{
1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
};
const u8 keccakf_piln[24] =
{
10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
};
/** /**
* 0x80 keccak, very special * 0x80 keccak, very special
*/ */

44
OpenCL/m05000_a3.cl
View File

@ -32,18 +32,6 @@ __constant u64 keccakf_rndc[24] =
0x8000000000008080, 0x0000000080000001, 0x8000000080008008 0x8000000000008080, 0x0000000080000001, 0x8000000080008008
}; };
__constant u8 keccakf_rotc[24] =
{
1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
};
__constant u8 keccakf_piln[24] =
{
10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
};
#ifndef KECCAK_ROUNDS #ifndef KECCAK_ROUNDS
#define KECCAK_ROUNDS 24 #define KECCAK_ROUNDS 24
#endif #endif
@ -91,6 +79,22 @@ static void m05000m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_le
const u32 gid = get_global_id (0); const u32 gid = get_global_id (0);
const u32 lid = get_local_id (0); const u32 lid = get_local_id (0);
/**
* constants
*/
const u8 keccakf_rotc[24] =
{
1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
};
const u8 keccakf_piln[24] =
{
10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
};
/** /**
* 0x80 keccak, very special * 0x80 keccak, very special
*/ */
@ -249,6 +253,22 @@ static void m05000s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_le
const u32 gid = get_global_id (0); const u32 gid = get_global_id (0);
const u32 lid = get_local_id (0); const u32 lid = get_local_id (0);
/**
* constants
*/
const u8 keccakf_rotc[24] =
{
1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
};
const u8 keccakf_piln[24] =
{
10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
};
/** /**
* 0x80 keccak, very special * 0x80 keccak, very special
*/ */