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

OpenCL Runtime: Workaround JiT compiler error on AMDGPU driver compiling WPA-EAPOL-PBKDF2 OpenCL kernel

Browse Source
This commit is contained in:
jsteube 2019-02-24 09:10:37 +01:00
parent 69c3ea2d27
commit 2c0c82c8af
2 changed files with 7 additions and 222 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

228
OpenCL/m02500-pure.cl
View File

@ -374,43 +374,7 @@ __kernel void m02500_aux1 (KERN_ATTR_TMPS_ESALT (wpa_pbkdf2_tmp_t, wpa_eapol_t))
sha1_hmac_init_64 (&ctx1, w0, w1, w2, w3);
ctx1.ipad.w0[0] = pke[ 0];
ctx1.ipad.w0[1] = pke[ 1];
ctx1.ipad.w0[2] = pke[ 2];
ctx1.ipad.w0[3] = pke[ 3];
ctx1.ipad.w1[0] = pke[ 4];
ctx1.ipad.w1[1] = pke[ 5];
ctx1.ipad.w1[2] = pke[ 6];
ctx1.ipad.w1[3] = pke[ 7];
ctx1.ipad.w2[0] = pke[ 8];
ctx1.ipad.w2[1] = pke[ 9];
ctx1.ipad.w2[2] = pke[10];
ctx1.ipad.w2[3] = pke[11];
ctx1.ipad.w3[0] = pke[12];
ctx1.ipad.w3[1] = pke[13];
ctx1.ipad.w3[2] = pke[14];
ctx1.ipad.w3[3] = pke[15];
sha1_transform (ctx1.ipad.w0, ctx1.ipad.w1, ctx1.ipad.w2, ctx1.ipad.w3, ctx1.ipad.h);
ctx1.ipad.w0[0] = pke[16];
ctx1.ipad.w0[1] = pke[17];
ctx1.ipad.w0[2] = pke[18];
ctx1.ipad.w0[3] = pke[19];
ctx1.ipad.w1[0] = pke[20];
ctx1.ipad.w1[1] = pke[21];
ctx1.ipad.w1[2] = pke[22];
ctx1.ipad.w1[3] = pke[23];
ctx1.ipad.w2[0] = pke[24];
ctx1.ipad.w2[1] = pke[25];
ctx1.ipad.w2[2] = pke[26];
ctx1.ipad.w2[3] = pke[27];
ctx1.ipad.w3[0] = pke[28];
ctx1.ipad.w3[1] = pke[29];
ctx1.ipad.w3[2] = pke[30];
ctx1.ipad.w3[3] = pke[31];
ctx1.ipad.len += 100;
sha1_hmac_update (&ctx1, pke, 100);
sha1_hmac_final (&ctx1);
@ -530,43 +494,7 @@ __kernel void m02500_aux1 (KERN_ATTR_TMPS_ESALT (wpa_pbkdf2_tmp_t, wpa_eapol_t))
sha1_hmac_init_64 (&ctx1, w0, w1, w2, w3);
ctx1.ipad.w0[0] = pke[ 0];
ctx1.ipad.w0[1] = pke[ 1];
ctx1.ipad.w0[2] = pke[ 2];
ctx1.ipad.w0[3] = pke[ 3];
ctx1.ipad.w1[0] = pke[ 4];
ctx1.ipad.w1[1] = pke[ 5];
ctx1.ipad.w1[2] = pke[ 6];
ctx1.ipad.w1[3] = pke[ 7];
ctx1.ipad.w2[0] = pke[ 8];
ctx1.ipad.w2[1] = pke[ 9];
ctx1.ipad.w2[2] = pke[10];
ctx1.ipad.w2[3] = pke[11];
ctx1.ipad.w3[0] = pke[12];
ctx1.ipad.w3[1] = pke[13];
ctx1.ipad.w3[2] = pke[14];
ctx1.ipad.w3[3] = pke[15];
sha1_transform (ctx1.ipad.w0, ctx1.ipad.w1, ctx1.ipad.w2, ctx1.ipad.w3, ctx1.ipad.h);
ctx1.ipad.w0[0] = pke[16];
ctx1.ipad.w0[1] = pke[17];
ctx1.ipad.w0[2] = pke[18];
ctx1.ipad.w0[3] = pke[19];
ctx1.ipad.w1[0] = pke[20];
ctx1.ipad.w1[1] = pke[21];
ctx1.ipad.w1[2] = pke[22];
ctx1.ipad.w1[3] = pke[23];
ctx1.ipad.w2[0] = pke[24];
ctx1.ipad.w2[1] = pke[25];
ctx1.ipad.w2[2] = pke[26];
ctx1.ipad.w2[3] = pke[27];
ctx1.ipad.w3[0] = pke[28];
ctx1.ipad.w3[1] = pke[29];
ctx1.ipad.w3[2] = pke[30];
ctx1.ipad.w3[3] = pke[31];
ctx1.ipad.len += 100;
sha1_hmac_update (&ctx1, pke, 100);
sha1_hmac_final (&ctx1);
@ -757,43 +685,7 @@ __kernel void m02500_aux2 (KERN_ATTR_TMPS_ESALT (wpa_pbkdf2_tmp_t, wpa_eapol_t))
sha1_hmac_init_64 (&ctx1, w0, w1, w2, w3);
ctx1.ipad.w0[0] = pke[ 0];
ctx1.ipad.w0[1] = pke[ 1];
ctx1.ipad.w0[2] = pke[ 2];
ctx1.ipad.w0[3] = pke[ 3];
ctx1.ipad.w1[0] = pke[ 4];
ctx1.ipad.w1[1] = pke[ 5];
ctx1.ipad.w1[2] = pke[ 6];
ctx1.ipad.w1[3] = pke[ 7];
ctx1.ipad.w2[0] = pke[ 8];
ctx1.ipad.w2[1] = pke[ 9];
ctx1.ipad.w2[2] = pke[10];
ctx1.ipad.w2[3] = pke[11];
ctx1.ipad.w3[0] = pke[12];
ctx1.ipad.w3[1] = pke[13];
ctx1.ipad.w3[2] = pke[14];
ctx1.ipad.w3[3] = pke[15];
sha1_transform (ctx1.ipad.w0, ctx1.ipad.w1, ctx1.ipad.w2, ctx1.ipad.w3, ctx1.ipad.h);
ctx1.ipad.w0[0] = pke[16];
ctx1.ipad.w0[1] = pke[17];
ctx1.ipad.w0[2] = pke[18];
ctx1.ipad.w0[3] = pke[19];
ctx1.ipad.w1[0] = pke[20];
ctx1.ipad.w1[1] = pke[21];
ctx1.ipad.w1[2] = pke[22];
ctx1.ipad.w1[3] = pke[23];
ctx1.ipad.w2[0] = pke[24];
ctx1.ipad.w2[1] = pke[25];
ctx1.ipad.w2[2] = pke[26];
ctx1.ipad.w2[3] = pke[27];
ctx1.ipad.w3[0] = pke[28];
ctx1.ipad.w3[1] = pke[29];
ctx1.ipad.w3[2] = pke[30];
ctx1.ipad.w3[3] = pke[31];
ctx1.ipad.len += 100;
sha1_hmac_update (&ctx1, pke, 100);
sha1_hmac_final (&ctx1);
@ -913,43 +805,7 @@ __kernel void m02500_aux2 (KERN_ATTR_TMPS_ESALT (wpa_pbkdf2_tmp_t, wpa_eapol_t))
sha1_hmac_init_64 (&ctx1, w0, w1, w2, w3);
ctx1.ipad.w0[0] = pke[ 0];
ctx1.ipad.w0[1] = pke[ 1];
ctx1.ipad.w0[2] = pke[ 2];
ctx1.ipad.w0[3] = pke[ 3];
ctx1.ipad.w1[0] = pke[ 4];
ctx1.ipad.w1[1] = pke[ 5];
ctx1.ipad.w1[2] = pke[ 6];
ctx1.ipad.w1[3] = pke[ 7];
ctx1.ipad.w2[0] = pke[ 8];
ctx1.ipad.w2[1] = pke[ 9];
ctx1.ipad.w2[2] = pke[10];
ctx1.ipad.w2[3] = pke[11];
ctx1.ipad.w3[0] = pke[12];
ctx1.ipad.w3[1] = pke[13];
ctx1.ipad.w3[2] = pke[14];
ctx1.ipad.w3[3] = pke[15];
sha1_transform (ctx1.ipad.w0, ctx1.ipad.w1, ctx1.ipad.w2, ctx1.ipad.w3, ctx1.ipad.h);
ctx1.ipad.w0[0] = pke[16];
ctx1.ipad.w0[1] = pke[17];
ctx1.ipad.w0[2] = pke[18];
ctx1.ipad.w0[3] = pke[19];
ctx1.ipad.w1[0] = pke[20];
ctx1.ipad.w1[1] = pke[21];
ctx1.ipad.w1[2] = pke[22];
ctx1.ipad.w1[3] = pke[23];
ctx1.ipad.w2[0] = pke[24];
ctx1.ipad.w2[1] = pke[25];
ctx1.ipad.w2[2] = pke[26];
ctx1.ipad.w2[3] = pke[27];
ctx1.ipad.w3[0] = pke[28];
ctx1.ipad.w3[1] = pke[29];
ctx1.ipad.w3[2] = pke[30];
ctx1.ipad.w3[3] = pke[31];
ctx1.ipad.len += 100;
sha1_hmac_update (&ctx1, pke, 100);
sha1_hmac_final (&ctx1);
@ -1192,43 +1048,7 @@ __kernel void m02500_aux3 (KERN_ATTR_TMPS_ESALT (wpa_pbkdf2_tmp_t, wpa_eapol_t))
sha256_hmac_init_64 (&ctx1, w0, w1, w2, w3);
ctx1.ipad.w0[0] = pke[ 0];
ctx1.ipad.w0[1] = pke[ 1];
ctx1.ipad.w0[2] = pke[ 2];
ctx1.ipad.w0[3] = pke[ 3];
ctx1.ipad.w1[0] = pke[ 4];
ctx1.ipad.w1[1] = pke[ 5];
ctx1.ipad.w1[2] = pke[ 6];
ctx1.ipad.w1[3] = pke[ 7];
ctx1.ipad.w2[0] = pke[ 8];
ctx1.ipad.w2[1] = pke[ 9];
ctx1.ipad.w2[2] = pke[10];
ctx1.ipad.w2[3] = pke[11];
ctx1.ipad.w3[0] = pke[12];
ctx1.ipad.w3[1] = pke[13];
ctx1.ipad.w3[2] = pke[14];
ctx1.ipad.w3[3] = pke[15];
sha256_transform (ctx1.ipad.w0, ctx1.ipad.w1, ctx1.ipad.w2, ctx1.ipad.w3, ctx1.ipad.h);
ctx1.ipad.w0[0] = pke[16];
ctx1.ipad.w0[1] = pke[17];
ctx1.ipad.w0[2] = pke[18];
ctx1.ipad.w0[3] = pke[19];
ctx1.ipad.w1[0] = pke[20];
ctx1.ipad.w1[1] = pke[21];
ctx1.ipad.w1[2] = pke[22];
ctx1.ipad.w1[3] = pke[23];
ctx1.ipad.w2[0] = pke[24];
ctx1.ipad.w2[1] = pke[25];
ctx1.ipad.w2[2] = pke[26];
ctx1.ipad.w2[3] = pke[27];
ctx1.ipad.w3[0] = pke[28];
ctx1.ipad.w3[1] = pke[29];
ctx1.ipad.w3[2] = pke[30];
ctx1.ipad.w3[3] = pke[31];
ctx1.ipad.len += 102;
sha256_hmac_update (&ctx1, pke, 102);
sha256_hmac_final (&ctx1);
@ -1379,43 +1199,7 @@ __kernel void m02500_aux3 (KERN_ATTR_TMPS_ESALT (wpa_pbkdf2_tmp_t, wpa_eapol_t))
sha256_hmac_init_64 (&ctx1, w0, w1, w2, w3);
ctx1.ipad.w0[0] = pke[ 0];
ctx1.ipad.w0[1] = pke[ 1];
ctx1.ipad.w0[2] = pke[ 2];
ctx1.ipad.w0[3] = pke[ 3];
ctx1.ipad.w1[0] = pke[ 4];
ctx1.ipad.w1[1] = pke[ 5];
ctx1.ipad.w1[2] = pke[ 6];
ctx1.ipad.w1[3] = pke[ 7];
ctx1.ipad.w2[0] = pke[ 8];
ctx1.ipad.w2[1] = pke[ 9];
ctx1.ipad.w2[2] = pke[10];
ctx1.ipad.w2[3] = pke[11];
ctx1.ipad.w3[0] = pke[12];
ctx1.ipad.w3[1] = pke[13];
ctx1.ipad.w3[2] = pke[14];
ctx1.ipad.w3[3] = pke[15];
sha256_transform (ctx1.ipad.w0, ctx1.ipad.w1, ctx1.ipad.w2, ctx1.ipad.w3, ctx1.ipad.h);
ctx1.ipad.w0[0] = pke[16];
ctx1.ipad.w0[1] = pke[17];
ctx1.ipad.w0[2] = pke[18];
ctx1.ipad.w0[3] = pke[19];
ctx1.ipad.w1[0] = pke[20];
ctx1.ipad.w1[1] = pke[21];
ctx1.ipad.w1[2] = pke[22];
ctx1.ipad.w1[3] = pke[23];
ctx1.ipad.w2[0] = pke[24];
ctx1.ipad.w2[1] = pke[25];
ctx1.ipad.w2[2] = pke[26];
ctx1.ipad.w2[3] = pke[27];
ctx1.ipad.w3[0] = pke[28];
ctx1.ipad.w3[1] = pke[29];
ctx1.ipad.w3[2] = pke[30];
ctx1.ipad.w3[3] = pke[31];
ctx1.ipad.len += 102;
sha256_hmac_update (&ctx1, pke, 102);
sha256_hmac_final (&ctx1);

1
docs/changes.txt
View File

@ -27,6 +27,7 @@
- Tuning Database: Updated hashcat.hctune with new models and refreshed vector width values
- Bitcoin Wallet: Be more user friendly by allowing a larger data range for ckey and public_key
- OpenCL Runtime: Workaround JiT compiler error on AMDGPU driver compiling WPA-EAPOL-PBKDF2 OpenCL kernel
- OpenCL Runtime: Improve ROCM detection and make sure to not confuse with recent AMDGPU drivers
##