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

Fixed invalid transfer from __constant to __local memory in -m 9100

Browse Source
This commit is contained in:
jsteube 2019-02-24 08:26:10 +01:00
parent be365acef8
commit 69c3ea2d27
1 changed files with 8 additions and 10 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

18
OpenCL/m09100-pure.cl
View File

@ -16,8 +16,7 @@
#define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "inc_comp_multi.cl"
// breaks if used with u8a on AMDGPU-PRO
__constant u8a lotus64_table[64] =
__constant u32 lotus64_table[64] =
{
'0', '1', '2', '3', '4', '5', '6', '7',
'8', '9', 'A', 'B', 'C', 'D', 'E', 'F',
@ -29,8 +28,7 @@ __constant u8a lotus64_table[64] =
'u', 'v', 'w', 'x', 'y', 'z', '+', '/',
};
// break if used with u8 on NVidia driver 378.x
__constant u8a lotus_magic_table[256] =
__constant u32 lotus_magic_table[256] =
{
0xbd, 0x56, 0xea, 0xf2, 0xa2, 0xf1, 0xac, 0x2a,
0xb0, 0x93, 0xd1, 0x9c, 0x1b, 0x33, 0xfd, 0xd0,
@ -70,7 +68,7 @@ __constant u8a lotus_magic_table[256] =
#define BOX1(S,i) (S)[(i)]
DECLSPEC void lotus_mix (u32 *in, const __local u8 *s_lotus_magic_table)
DECLSPEC void lotus_mix (u32 *in, const __local u32 *s_lotus_magic_table)
{
u8 p = 0;
@ -93,7 +91,7 @@ DECLSPEC void lotus_mix (u32 *in, const __local u8 *s_lotus_magic_table)
}
}
DECLSPEC void lotus_transform_password (const u32 *in, u32 *out, const __local u8 *s_lotus_magic_table)
DECLSPEC void lotus_transform_password (const u32 *in, u32 *out, const __local u32 *s_lotus_magic_table)
{
u8 t = (u8) (out[3] >> 24);
@ -190,7 +188,7 @@ DECLSPEC void pad (u32 *w, const u32 len)
}
}
DECLSPEC void mdtransform_norecalc (u32 *state, const u32 *block, const __local u8 *s_lotus_magic_table)
DECLSPEC void mdtransform_norecalc (u32 *state, const u32 *block, const __local u32 *s_lotus_magic_table)
{
u32 x[12];
@ -215,14 +213,14 @@ DECLSPEC void mdtransform_norecalc (u32 *state, const u32 *block, const __local
state[3] = x[3];
}
DECLSPEC void mdtransform (u32 *state, u32 *checksum, const u32 *block, const __local u8 *s_lotus_magic_table)
DECLSPEC void mdtransform (u32 *state, u32 *checksum, const u32 *block, const __local u32 *s_lotus_magic_table)
{
mdtransform_norecalc (state, block, s_lotus_magic_table);
lotus_transform_password (block, checksum, s_lotus_magic_table);
}
DECLSPEC void domino_big_md (const u32 *saved_key, const u32 size, u32 *state, const __local u8 *s_lotus_magic_table)
DECLSPEC void domino_big_md (const u32 *saved_key, const u32 size, u32 *state, const __local u32 *s_lotus_magic_table)
{
u32 checksum[4];
@ -400,7 +398,7 @@ __kernel void m09100_init (KERN_ATTR_TMPS (lotus8_tmp_t))
* sbox
*/
__local u8 s_lotus_magic_table[256];
__local u32 s_lotus_magic_table[256];
for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
{