Revert "Workaround buggy NV OpenCL runtime"

This reverts commit 90eba9cd2e1994b6eb01a858e4603936e2c50e83.

diff --git a/OpenCL/m05500_a0.cl b/OpenCL/m05500_a0.cl
index fce3625..64d1a40 100644 (file)
--- a/OpenCL/m05500_a0.cl
+++ b/OpenCL/m05500_a0.cl
@@ -673,10 +673,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_m04 (__glo
      * DES2
      */
 
-    volatile const u32 bc = (b >> 24) | (c << 8);
-    volatile const u32 cd = (c >> 24) | (d << 8);
-
-    transform_netntlmv1_key (bc, cd, key);
+    transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key);
 
     _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
      */
 
-    volatile const u32 bc = (b >> 24) | (c << 8);
-    volatile const u32 cd = (c >> 24) | (d << 8);
-
-    transform_netntlmv1_key (bc, cd, key);
+    /*
+    transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), 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

diff --git a/OpenCL/m05500_a1.cl b/OpenCL/m05500_a1.cl
index 81d7650..4d44506 100644 (file)
--- a/OpenCL/m05500_a1.cl
+++ b/OpenCL/m05500_a1.cl
@@ -725,10 +725,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_m04 (__glo
      * DES2
      */
 
-    volatile const u32 bc = (b >> 24) | (c << 8);
-    volatile const u32 cd = (c >> 24) | (d << 8);
-
-    transform_netntlmv1_key (bc, cd, key);
+    transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key);
 
     _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
      */
 
-    volatile const u32 bc = (b >> 24) | (c << 8);
-    volatile const u32 cd = (c >> 24) | (d << 8);
-
-    transform_netntlmv1_key (bc, cd, key);
+    /*
+    transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), 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

diff --git a/OpenCL/m05500_a3.cl b/OpenCL/m05500_a3.cl
index 4b08e14..3a507f4 100644 (file)
--- a/OpenCL/m05500_a3.cl
+++ b/OpenCL/m05500_a3.cl
@@ -602,7 +602,6 @@ 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);
@@ -610,12 +609,6 @@ 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
@@ -778,16 +771,20 @@ static void m05500s (__local u32 s_SPtrans[8][64], __local u32 s_skb[8][64], u32
      * DES2
      */
 
-    volatile const u32 bc = (b >> 24) | (c << 8);
-    volatile const u32 cd = (c >> 24) | (d << 8);
-
-    transform_netntlmv1_key (bc, cd, key);
+    /*
+    transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), 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