arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2025-05-29 04:08:48 +00:00
Code Issues Releases Wiki Activity

Revert "Workaround buggy NV OpenCL runtime"

Browse Source 
This reverts commit 90eba9cd2e.
This commit is contained in:
jsteube 2015-12-24 12:05:43 +01:00
parent 90eba9cd2e
commit 0522ca9cbe
3 changed files with 26 additions and 27 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

17
OpenCL/m05500_a0.cl
View File

@ -673,10 +673,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_m04 (__glo
* DES2 * DES2
*/ */
volatile const u32 bc = (b >> 24) | (c << 8); transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key);
volatile const u32 cd = (c >> 24) | (d << 8);
transform_netntlmv1_key (bc, cd, key);
_des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb); _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
@ -922,16 +919,20 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_s04 (__glo
* DES2 * DES2
*/ */
volatile const u32 bc = (b >> 24) | (c << 8); /*
volatile const u32 cd = (c >> 24) | (d << 8); transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key);
transform_netntlmv1_key (bc, cd, key);
_des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb); _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
u32 iv2[2]; u32 iv2[2];
_des_crypt_encrypt (iv2, data, Kc, Kd, s_SPtrans); _des_crypt_encrypt (iv2, data, Kc, Kd, s_SPtrans);
*/
u32 iv2[2];
iv2[0] = search[2];
iv2[1] = search[3];
/** /**
* compare * compare

17
OpenCL/m05500_a1.cl
View File

@ -725,10 +725,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_m04 (__glo
* DES2 * DES2
*/ */
volatile const u32 bc = (b >> 24) | (c << 8); transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key);
volatile const u32 cd = (c >> 24) | (d << 8);
transform_netntlmv1_key (bc, cd, key);
_des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb); _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
@ -1028,16 +1025,20 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_s04 (__glo
* DES2 * DES2
*/ */
volatile const u32 bc = (b >> 24) | (c << 8); /*
volatile const u32 cd = (c >> 24) | (d << 8); transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key);
transform_netntlmv1_key (bc, cd, key);
_des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb); _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
u32 iv2[2]; u32 iv2[2];
_des_crypt_encrypt (iv2, data, Kc, Kd, s_SPtrans); _des_crypt_encrypt (iv2, data, Kc, Kd, s_SPtrans);
*/
u32 iv2[2];
iv2[0] = search[2];
iv2[1] = search[3];
/** /**
* compare * compare

19
OpenCL/m05500_a3.cl
View File

@ -602,7 +602,6 @@ static void m05500m (__local u32 s_SPtrans[8][64], __local u32 s_skb[8][64], u32
* DES2 * DES2
*/ */
/*
transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key); transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key);
_des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb); _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
@ -610,12 +609,6 @@ static void m05500m (__local u32 s_SPtrans[8][64], __local u32 s_skb[8][64], u32
u32 iv2[2]; u32 iv2[2];
_des_crypt_encrypt (iv2, data, Kc, Kd, s_SPtrans); _des_crypt_encrypt (iv2, data, Kc, Kd, s_SPtrans);
*/
u32 iv2[2];
iv2[0] = search[2];
iv2[1] = search[3];
/** /**
* compare * compare
@ -778,16 +771,20 @@ static void m05500s (__local u32 s_SPtrans[8][64], __local u32 s_skb[8][64], u32
* DES2 * DES2
*/ */
volatile const u32 bc = (b >> 24) | (c << 8); /*
volatile const u32 cd = (c >> 24) | (d << 8); transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key);
transform_netntlmv1_key (bc, cd, key);
_des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb); _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
u32 iv2[2]; u32 iv2[2];
_des_crypt_encrypt (iv2, data, Kc, Kd, s_SPtrans); _des_crypt_encrypt (iv2, data, Kc, Kd, s_SPtrans);
*/
u32 iv2[2];
iv2[0] = search[2];
iv2[1] = search[3];
/** /**
* compare * compare