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

Merge pull request #617 from matrix/truecrypt_osx

Browse Source 
Fix Truecrypt OSX OpenCL kernels
This commit is contained in:
Jens Steube 2016-11-26 23:17:16 +01:00 committed by GitHub
commit a1374f1d0a
3 changed files with 28 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

12
OpenCL/m06213.cl
View File

@ -707,7 +707,11 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul
ukey3[6] = tmps[gid].out[22]; ukey3[6] = tmps[gid].out[22];
ukey3[7] = tmps[gid].out[23]; ukey3[7] = tmps[gid].out[23];
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey4[8];
#else
u32 ukey4[8]; u32 ukey4[8];
#endif
ukey4[0] = tmps[gid].out[24]; ukey4[0] = tmps[gid].out[24];
ukey4[1] = tmps[gid].out[25]; ukey4[1] = tmps[gid].out[25];
@ -763,7 +767,11 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul
} }
} }
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey5[8];
#else
u32 ukey5[8]; u32 ukey5[8];
#endif
ukey5[0] = tmps[gid].out[32]; ukey5[0] = tmps[gid].out[32];
ukey5[1] = tmps[gid].out[33]; ukey5[1] = tmps[gid].out[33];
@ -774,7 +782,11 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul
ukey5[6] = tmps[gid].out[38]; ukey5[6] = tmps[gid].out[38];
ukey5[7] = tmps[gid].out[39]; ukey5[7] = tmps[gid].out[39];
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey6[8];
#else
u32 ukey6[8]; u32 ukey6[8];
#endif
ukey6[0] = tmps[gid].out[40]; ukey6[0] = tmps[gid].out[40];
ukey6[1] = tmps[gid].out[41]; ukey6[1] = tmps[gid].out[41];

4
OpenCL/m06223.cl
View File

@ -615,7 +615,11 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul
ukey3[6] = swap32 (h32_from_64 (tmps[gid].out[11])); ukey3[6] = swap32 (h32_from_64 (tmps[gid].out[11]));
ukey3[7] = swap32 (l32_from_64 (tmps[gid].out[11])); ukey3[7] = swap32 (l32_from_64 (tmps[gid].out[11]));
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey4[8];
#else
u32 ukey4[8]; u32 ukey4[8];
#endif
ukey4[0] = swap32 (h32_from_64 (tmps[gid].out[12])); ukey4[0] = swap32 (h32_from_64 (tmps[gid].out[12]));
ukey4[1] = swap32 (l32_from_64 (tmps[gid].out[12])); ukey4[1] = swap32 (l32_from_64 (tmps[gid].out[12]));

12
OpenCL/m06233.cl
View File

@ -2019,7 +2019,11 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul
ukey3[6] = swap32 (tmps[gid].out[22]); ukey3[6] = swap32 (tmps[gid].out[22]);
ukey3[7] = swap32 (tmps[gid].out[23]); ukey3[7] = swap32 (tmps[gid].out[23]);
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey4[8];
#else
u32 ukey4[8]; u32 ukey4[8];
#endif
ukey4[0] = swap32 (tmps[gid].out[24]); ukey4[0] = swap32 (tmps[gid].out[24]);
ukey4[1] = swap32 (tmps[gid].out[25]); ukey4[1] = swap32 (tmps[gid].out[25]);
@ -2075,7 +2079,11 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul
} }
} }
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey5[8];
#else
u32 ukey5[8]; u32 ukey5[8];
#endif
ukey5[0] = swap32 (tmps[gid].out[32]); ukey5[0] = swap32 (tmps[gid].out[32]);
ukey5[1] = swap32 (tmps[gid].out[33]); ukey5[1] = swap32 (tmps[gid].out[33]);
@ -2086,7 +2094,11 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul
ukey5[6] = swap32 (tmps[gid].out[38]); ukey5[6] = swap32 (tmps[gid].out[38]);
ukey5[7] = swap32 (tmps[gid].out[39]); ukey5[7] = swap32 (tmps[gid].out[39]);
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey6[8];
#else
u32 ukey6[8]; u32 ukey6[8];
#endif
ukey6[0] = swap32 (tmps[gid].out[40]); ukey6[0] = swap32 (tmps[gid].out[40]);
ukey6[1] = swap32 (tmps[gid].out[41]); ukey6[1] = swap32 (tmps[gid].out[41]);