From 2c0c82c8affd800e6f66b0d61b959534f1bceb43 Mon Sep 17 00:00:00 2001
From: jsteube <jens.steube@gmail.com>
Date: Sun, 24 Feb 2019 09:10:37 +0100
Subject: [PATCH] OpenCL Runtime: Workaround JiT compiler error on AMDGPU
 driver compiling WPA-EAPOL-PBKDF2 OpenCL kernel

---
 OpenCL/m02500-pure.cl | 228 ++----------------------------------------
 docs/changes.txt      |   1 +
 2 files changed, 7 insertions(+), 222 deletions(-)

diff --git a/OpenCL/m02500-pure.cl b/OpenCL/m02500-pure.cl
index 5d66edf03..5241a6bd9 100644
--- a/OpenCL/m02500-pure.cl
+++ b/OpenCL/m02500-pure.cl
@@ -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);
 
diff --git a/docs/changes.txt b/docs/changes.txt
index afd99aa07..a03261341 100644
--- a/docs/changes.txt
+++ b/docs/changes.txt
@@ -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
 
 ##