arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2025-06-20 23:18:48 +00:00
Code Issues Releases Wiki Activity

Get rid of volatile in TrueCrypt kernels

Browse Source
This commit is contained in:
jsteube 2017-09-08 19:47:56 +02:00
parent 16e33b20fc
commit 55f653f374
6 changed files with 2 additions and 122 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

16
OpenCL/m06212.cl
View File

@ -331,11 +331,7 @@ __kernel void m06212_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (gid >= gid_max) return; if (gid >= gid_max) return;
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey1[8];
#else
u32 ukey1[8]; u32 ukey1[8];
#endif
ukey1[0] = tmps[gid].out[ 0]; ukey1[0] = tmps[gid].out[ 0];
ukey1[1] = tmps[gid].out[ 1]; ukey1[1] = tmps[gid].out[ 1];
@ -346,11 +342,7 @@ __kernel void m06212_comp (__global pw_t *pws, __global const kernel_rule_t *rul
ukey1[6] = tmps[gid].out[ 6]; ukey1[6] = tmps[gid].out[ 6];
ukey1[7] = tmps[gid].out[ 7]; ukey1[7] = tmps[gid].out[ 7];
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey2[8];
#else
u32 ukey2[8]; u32 ukey2[8];
#endif
ukey2[0] = tmps[gid].out[ 8]; ukey2[0] = tmps[gid].out[ 8];
ukey2[1] = tmps[gid].out[ 9]; ukey2[1] = tmps[gid].out[ 9];
@ -385,11 +377,7 @@ __kernel void m06212_comp (__global pw_t *pws, __global const kernel_rule_t *rul
} }
} }
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey3[8];
#else
u32 ukey3[8]; u32 ukey3[8];
#endif
ukey3[0] = tmps[gid].out[16]; ukey3[0] = tmps[gid].out[16];
ukey3[1] = tmps[gid].out[17]; ukey3[1] = tmps[gid].out[17];
@ -400,11 +388,7 @@ __kernel void m06212_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];

24
OpenCL/m06213.cl
View File

@ -331,11 +331,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (gid >= gid_max) return; if (gid >= gid_max) return;
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey1[8];
#else
u32 ukey1[8]; u32 ukey1[8];
#endif
ukey1[0] = tmps[gid].out[ 0]; ukey1[0] = tmps[gid].out[ 0];
ukey1[1] = tmps[gid].out[ 1]; ukey1[1] = tmps[gid].out[ 1];
@ -346,11 +342,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul
ukey1[6] = tmps[gid].out[ 6]; ukey1[6] = tmps[gid].out[ 6];
ukey1[7] = tmps[gid].out[ 7]; ukey1[7] = tmps[gid].out[ 7];
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey2[8];
#else
u32 ukey2[8]; u32 ukey2[8];
#endif
ukey2[0] = tmps[gid].out[ 8]; ukey2[0] = tmps[gid].out[ 8];
ukey2[1] = tmps[gid].out[ 9]; ukey2[1] = tmps[gid].out[ 9];
@ -385,11 +377,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul
} }
} }
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey3[8];
#else
u32 ukey3[8]; u32 ukey3[8];
#endif
ukey3[0] = tmps[gid].out[16]; ukey3[0] = tmps[gid].out[16];
ukey3[1] = tmps[gid].out[17]; ukey3[1] = tmps[gid].out[17];
@ -400,11 +388,7 @@ __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];
@ -439,11 +423,7 @@ __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];
@ -454,11 +434,7 @@ __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];

16
OpenCL/m06222.cl
View File

@ -475,11 +475,7 @@ __kernel void m06222_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (gid >= gid_max) return; if (gid >= gid_max) return;
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey1[8];
#else
u32 ukey1[8]; u32 ukey1[8];
#endif
ukey1[0] = swap32_S (h32_from_64_S (tmps[gid].out[0])); ukey1[0] = swap32_S (h32_from_64_S (tmps[gid].out[0]));
ukey1[1] = swap32_S (l32_from_64_S (tmps[gid].out[0])); ukey1[1] = swap32_S (l32_from_64_S (tmps[gid].out[0]));
@ -490,11 +486,7 @@ __kernel void m06222_comp (__global pw_t *pws, __global const kernel_rule_t *rul
ukey1[6] = swap32_S (h32_from_64_S (tmps[gid].out[3])); ukey1[6] = swap32_S (h32_from_64_S (tmps[gid].out[3]));
ukey1[7] = swap32_S (l32_from_64_S (tmps[gid].out[3])); ukey1[7] = swap32_S (l32_from_64_S (tmps[gid].out[3]));
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey2[8];
#else
u32 ukey2[8]; u32 ukey2[8];
#endif
ukey2[0] = swap32_S (h32_from_64_S (tmps[gid].out[4])); ukey2[0] = swap32_S (h32_from_64_S (tmps[gid].out[4]));
ukey2[1] = swap32_S (l32_from_64_S (tmps[gid].out[4])); ukey2[1] = swap32_S (l32_from_64_S (tmps[gid].out[4]));
@ -529,11 +521,7 @@ __kernel void m06222_comp (__global pw_t *pws, __global const kernel_rule_t *rul
} }
} }
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey3[8];
#else
u32 ukey3[8]; u32 ukey3[8];
#endif
ukey3[0] = swap32_S (h32_from_64_S (tmps[gid].out[ 8])); ukey3[0] = swap32_S (h32_from_64_S (tmps[gid].out[ 8]));
ukey3[1] = swap32_S (l32_from_64_S (tmps[gid].out[ 8])); ukey3[1] = swap32_S (l32_from_64_S (tmps[gid].out[ 8]));
@ -544,11 +532,7 @@ __kernel void m06222_comp (__global pw_t *pws, __global const kernel_rule_t *rul
ukey3[6] = swap32_S (h32_from_64_S (tmps[gid].out[11])); ukey3[6] = swap32_S (h32_from_64_S (tmps[gid].out[11]));
ukey3[7] = swap32_S (l32_from_64_S (tmps[gid].out[11])); ukey3[7] = swap32_S (l32_from_64_S (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_S (h32_from_64_S (tmps[gid].out[12])); ukey4[0] = swap32_S (h32_from_64_S (tmps[gid].out[12]));
ukey4[1] = swap32_S (l32_from_64_S (tmps[gid].out[12])); ukey4[1] = swap32_S (l32_from_64_S (tmps[gid].out[12]));

28
OpenCL/m06223.cl
View File

@ -475,11 +475,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (gid >= gid_max) return; if (gid >= gid_max) return;
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey1[8];
#else
u32 ukey1[8]; u32 ukey1[8];
#endif
ukey1[0] = swap32_S (h32_from_64_S (tmps[gid].out[0])); ukey1[0] = swap32_S (h32_from_64_S (tmps[gid].out[0]));
ukey1[1] = swap32_S (l32_from_64_S (tmps[gid].out[0])); ukey1[1] = swap32_S (l32_from_64_S (tmps[gid].out[0]));
@ -490,11 +486,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul
ukey1[6] = swap32_S (h32_from_64_S (tmps[gid].out[3])); ukey1[6] = swap32_S (h32_from_64_S (tmps[gid].out[3]));
ukey1[7] = swap32_S (l32_from_64_S (tmps[gid].out[3])); ukey1[7] = swap32_S (l32_from_64_S (tmps[gid].out[3]));
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey2[8];
#else
u32 ukey2[8]; u32 ukey2[8];
#endif
ukey2[0] = swap32_S (h32_from_64_S (tmps[gid].out[4])); ukey2[0] = swap32_S (h32_from_64_S (tmps[gid].out[4]));
ukey2[1] = swap32_S (l32_from_64_S (tmps[gid].out[4])); ukey2[1] = swap32_S (l32_from_64_S (tmps[gid].out[4]));
@ -529,11 +521,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul
} }
} }
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey3[8];
#else
u32 ukey3[8]; u32 ukey3[8];
#endif
ukey3[0] = swap32_S (h32_from_64_S (tmps[gid].out[ 8])); ukey3[0] = swap32_S (h32_from_64_S (tmps[gid].out[ 8]));
ukey3[1] = swap32_S (l32_from_64_S (tmps[gid].out[ 8])); ukey3[1] = swap32_S (l32_from_64_S (tmps[gid].out[ 8]));
@ -544,11 +532,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul
ukey3[6] = swap32_S (h32_from_64_S (tmps[gid].out[11])); ukey3[6] = swap32_S (h32_from_64_S (tmps[gid].out[11]));
ukey3[7] = swap32_S (l32_from_64_S (tmps[gid].out[11])); ukey3[7] = swap32_S (l32_from_64_S (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_S (h32_from_64_S (tmps[gid].out[12])); ukey4[0] = swap32_S (h32_from_64_S (tmps[gid].out[12]));
ukey4[1] = swap32_S (l32_from_64_S (tmps[gid].out[12])); ukey4[1] = swap32_S (l32_from_64_S (tmps[gid].out[12]));
@ -583,11 +567,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul
} }
} }
#if defined (IS_APPLE) && defined (IS_GPU) u32 ukey5[8];
volatile u32 ukey5[8];
#else
volatile u32 ukey5[8];
#endif
ukey5[0] = swap32_S (h32_from_64_S (tmps[gid].out[16])); ukey5[0] = swap32_S (h32_from_64_S (tmps[gid].out[16]));
ukey5[1] = swap32_S (l32_from_64_S (tmps[gid].out[16])); ukey5[1] = swap32_S (l32_from_64_S (tmps[gid].out[16]));
@ -598,11 +578,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul
ukey5[6] = swap32_S (h32_from_64_S (tmps[gid].out[19])); ukey5[6] = swap32_S (h32_from_64_S (tmps[gid].out[19]));
ukey5[7] = swap32_S (l32_from_64_S (tmps[gid].out[19])); ukey5[7] = swap32_S (l32_from_64_S (tmps[gid].out[19]));
#if defined (IS_APPLE) && defined (IS_GPU) u32 ukey6[8];
volatile u32 ukey6[8];
#else
volatile u32 ukey6[8];
#endif
ukey6[0] = swap32_S (h32_from_64_S (tmps[gid].out[20])); ukey6[0] = swap32_S (h32_from_64_S (tmps[gid].out[20]));
ukey6[1] = swap32_S (l32_from_64_S (tmps[gid].out[20])); ukey6[1] = swap32_S (l32_from_64_S (tmps[gid].out[20]));

16
OpenCL/m06232.cl
View File

@ -615,11 +615,7 @@ __kernel void m06232_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (gid >= gid_max) return; if (gid >= gid_max) return;
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey1[8];
#else
u32 ukey1[8]; u32 ukey1[8];
#endif
ukey1[0] = swap32_S (tmps[gid].out[ 0]); ukey1[0] = swap32_S (tmps[gid].out[ 0]);
ukey1[1] = swap32_S (tmps[gid].out[ 1]); ukey1[1] = swap32_S (tmps[gid].out[ 1]);
@ -630,11 +626,7 @@ __kernel void m06232_comp (__global pw_t *pws, __global const kernel_rule_t *rul
ukey1[6] = swap32_S (tmps[gid].out[ 6]); ukey1[6] = swap32_S (tmps[gid].out[ 6]);
ukey1[7] = swap32_S (tmps[gid].out[ 7]); ukey1[7] = swap32_S (tmps[gid].out[ 7]);
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey2[8];
#else
u32 ukey2[8]; u32 ukey2[8];
#endif
ukey2[0] = swap32_S (tmps[gid].out[ 8]); ukey2[0] = swap32_S (tmps[gid].out[ 8]);
ukey2[1] = swap32_S (tmps[gid].out[ 9]); ukey2[1] = swap32_S (tmps[gid].out[ 9]);
@ -669,11 +661,7 @@ __kernel void m06232_comp (__global pw_t *pws, __global const kernel_rule_t *rul
} }
} }
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey3[8];
#else
u32 ukey3[8]; u32 ukey3[8];
#endif
ukey3[0] = swap32_S (tmps[gid].out[16]); ukey3[0] = swap32_S (tmps[gid].out[16]);
ukey3[1] = swap32_S (tmps[gid].out[17]); ukey3[1] = swap32_S (tmps[gid].out[17]);
@ -684,11 +672,7 @@ __kernel void m06232_comp (__global pw_t *pws, __global const kernel_rule_t *rul
ukey3[6] = swap32_S (tmps[gid].out[22]); ukey3[6] = swap32_S (tmps[gid].out[22]);
ukey3[7] = swap32_S (tmps[gid].out[23]); ukey3[7] = swap32_S (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_S (tmps[gid].out[24]); ukey4[0] = swap32_S (tmps[gid].out[24]);
ukey4[1] = swap32_S (tmps[gid].out[25]); ukey4[1] = swap32_S (tmps[gid].out[25]);

24
OpenCL/m06233.cl
View File

@ -615,11 +615,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul
if (gid >= gid_max) return; if (gid >= gid_max) return;
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey1[8];
#else
u32 ukey1[8]; u32 ukey1[8];
#endif
ukey1[0] = swap32_S (tmps[gid].out[ 0]); ukey1[0] = swap32_S (tmps[gid].out[ 0]);
ukey1[1] = swap32_S (tmps[gid].out[ 1]); ukey1[1] = swap32_S (tmps[gid].out[ 1]);
@ -630,11 +626,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul
ukey1[6] = swap32_S (tmps[gid].out[ 6]); ukey1[6] = swap32_S (tmps[gid].out[ 6]);
ukey1[7] = swap32_S (tmps[gid].out[ 7]); ukey1[7] = swap32_S (tmps[gid].out[ 7]);
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey2[8];
#else
u32 ukey2[8]; u32 ukey2[8];
#endif
ukey2[0] = swap32_S (tmps[gid].out[ 8]); ukey2[0] = swap32_S (tmps[gid].out[ 8]);
ukey2[1] = swap32_S (tmps[gid].out[ 9]); ukey2[1] = swap32_S (tmps[gid].out[ 9]);
@ -669,11 +661,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul
} }
} }
#if defined (IS_APPLE) && defined (IS_GPU)
volatile u32 ukey3[8];
#else
u32 ukey3[8]; u32 ukey3[8];
#endif
ukey3[0] = swap32_S (tmps[gid].out[16]); ukey3[0] = swap32_S (tmps[gid].out[16]);
ukey3[1] = swap32_S (tmps[gid].out[17]); ukey3[1] = swap32_S (tmps[gid].out[17]);
@ -684,11 +672,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul
ukey3[6] = swap32_S (tmps[gid].out[22]); ukey3[6] = swap32_S (tmps[gid].out[22]);
ukey3[7] = swap32_S (tmps[gid].out[23]); ukey3[7] = swap32_S (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_S (tmps[gid].out[24]); ukey4[0] = swap32_S (tmps[gid].out[24]);
ukey4[1] = swap32_S (tmps[gid].out[25]); ukey4[1] = swap32_S (tmps[gid].out[25]);
@ -723,11 +707,7 @@ __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_S (tmps[gid].out[32]); ukey5[0] = swap32_S (tmps[gid].out[32]);
ukey5[1] = swap32_S (tmps[gid].out[33]); ukey5[1] = swap32_S (tmps[gid].out[33]);
@ -738,11 +718,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul
ukey5[6] = swap32_S (tmps[gid].out[38]); ukey5[6] = swap32_S (tmps[gid].out[38]);
ukey5[7] = swap32_S (tmps[gid].out[39]); ukey5[7] = swap32_S (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_S (tmps[gid].out[40]); ukey6[0] = swap32_S (tmps[gid].out[40]);
ukey6[1] = swap32_S (tmps[gid].out[41]); ukey6[1] = swap32_S (tmps[gid].out[41]);