1
0
mirror of https://github.com/hashcat/hashcat.git synced 2024-12-22 14:48:12 +00:00

Workaround buggy NV OpenCL runtime

This commit is contained in:
jsteube 2015-12-24 12:01:39 +01:00
parent 3dc8d526db
commit 90eba9cd2e
3 changed files with 27 additions and 26 deletions

View File

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

View File

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

View File

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